From 556c4e154a958cf575d6563095ad60935d1e3df5 Mon Sep 17 00:00:00 2001 From: pgovyadi Date: Mon, 29 Jun 2015 20:09:20 -0500 Subject: [PATCH] Changed the handling of the transformation matrix in gl_spider.h, added functions for handling the update implementation (compatible with vectors). Added comments and fixes to cost.h, minor changes to matrix.h, and vect.h in order to facilitate moving into 4D vector notation. --- stim/cuda/cost.h | 47 ++++++++++++++++++++++++----------------------- stim/gl/gl_spider.h | 70 ++++++++++++++++++++++++++++++++++++++++------------------------------ stim/math/vector.h | 1 - 3 files changed, 64 insertions(+), 54 deletions(-) diff --git a/stim/cuda/cost.h b/stim/cuda/cost.h index 6c1985a..4da0c80 100644 --- a/stim/cuda/cost.h +++ b/stim/cuda/cost.h @@ -6,13 +6,9 @@ #include "../visualization/colormap.h" #include -//#define -//#define DIM_Y 10890 -//#define DIM_X 20 + +///Cost function that works with the gl-spider class to find index of the item with min-cost. typedef unsigned char uchar; -//surface texOut; ///// maybe just do a normal array instead of a surface. - //we may not need a surface at all. -//texture texTemplate texture texIn; float *result; float* v_dif; @@ -28,7 +24,9 @@ inline void checkCUDAerrors(const char *msg) } } - +///Finds the sum of all the pixes in a gives template element. +///Returns the abosolute value. +///@param *diff, a pointer to the memory block that holds the pixel-differences. float get_sum(float *diff) { @@ -43,6 +41,9 @@ float get_sum(float *diff) return out; } +///A virtual representation of a uniform template. +///Returns the value of the template pixel. +///@param x, location of a pixel. __device__ float Template(int x) { if(x < 20/6 || x > 20*5/6 || (x > 20*2/6 && x < 20*4/6)){ @@ -53,6 +54,9 @@ __device__ float Template(int x) } +///Find the difference of the given set of samples and the template +///using cuda acceleration. +///@param *result, a pointer to the memory that stores the result. __global__ void get_diff (float *result) { @@ -61,25 +65,17 @@ void get_diff (float *result) int y = threadIdx.y + blockIdx.y * blockDim.y; int idx = y*20+x; - //float valIn = tex2D(texIn, x, y); float valIn = tex2D(texIn, x, y)/255.0; - //int idx = x*DIM_X+y; - - //uchar4 color = tex2D(texIn, x, y); - //float3 tempcolor = make_float3(color.x, color.y, color.z); - //float valIn = tempcolor.x + tempcolor.y + tempcolor.z; - //float valIn = tex2D(texIn, x, y); float valTemp = Template(x); -// result[idx] = (float) x/(blockDim.x*gridDim.x);//abs(x); - result[idx] = abs(valIn); -// #if __CUDA_ARCH__>=200 -// printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]); -// #endif + result[idx] = abs(valIn-valTemp); } - +///Initialization function, allocates the memory and passes the necessary +///handles from OpenGL and Cuda. +///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture +///@param DIM_Y, integer controlling how much memory to allocate. void initArray(cudaGraphicsResource_t src, int DIM_Y) { //cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc (); @@ -105,7 +101,9 @@ void initArray(cudaGraphicsResource_t src, int DIM_Y) // cudaBindTextureToArray(texIn, ptr, &channelDesc) // ); } - +///Deinit function that frees the memery used and releases the texture resource +///back to OpenGL. +///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture void cleanUP(cudaGraphicsResource_t src) { HANDLE_ERROR( @@ -121,14 +119,17 @@ void cleanUP(cudaGraphicsResource_t src) cudaFree(v_dif) ); } - +///External access-point to the cuda function +///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture +///@param DIM_Y, the number of samples in the template. +///@inter temporary paramenter that tracks the number of times cost.h was called. extern "C" int get_cost(cudaGraphicsResource_t src, int inter, int DIM_Y) { float output[DIM_Y]; float mini = 10000000000000000.0; int idx; - stringstream name; + stringstream name; //for debugging initArray(src, DIM_Y*10); dim3 grid(20, DIM_Y*10); dim3 block(1, 1); diff --git a/stim/gl/gl_spider.h b/stim/gl/gl_spider.h index ab1ae75..cbe2b3b 100644 --- a/stim/gl/gl_spider.h +++ b/stim/gl/gl_spider.h @@ -7,13 +7,14 @@ #include #include #include -#include "../gl/gl_texture.h" -#include "../visualization/camera.h" -#include "./error.h" -#include "../math/vector.h" -#include "../math/rect.h" -#include "../cuda/cost.h" -#include "../cuda/glbind.h" +#include "stim/gl/gl_texture.h" +#include "stim/visualization/camera.h" +#include "stim/gl/error.h" +#include "stim/math/vector.h" +#include "stim/math/rect.h" +#include "stim/math/matrix.h" +#include "stim/cuda/cost.h" +#include "stim/cuda/glbind.h" #include #include @@ -42,7 +43,7 @@ class gl_spider : public virtual gl_texture std::vector > dirVectors; std::vector > posVectors; std::vector > magVectors; - double currentTransform[16]; + stim::matrix currentTransform; using gl_texture::texID; using gl_texture::S; using gl_texture::R; @@ -62,9 +63,16 @@ class gl_spider : public virtual gl_texture { genTemplate(dirVectors, 0); int best = getCost(); - stim::vec next = dirVectors[best]; - // next[0] = next[0]* - + stim::vec next; + next[0] = dirVectors[best][0]*512.0*0.6; + next[1] = dirVectors[best][1]*512.0*0.6; + next[2] = dirVectors[best][2]*512.0*0.6; + next[3] = 1; + next = (currentTransform*next).norm(); + setPosition( position[0]+next[0]*magnitude[0]/4, + position[1]+next[1]*magnitude[0]/4, + position[2]+next[2]*magnitude[0]/4); + setDirection(next[0], next[1], next[2]); } /// Method for finding the best direction for the spider. @@ -419,23 +427,19 @@ class gl_spider : public virtual gl_texture stim::vec rot = getRotation(direction); glMatrixMode(GL_TEXTURE); glLoadIdentity(); - glScalef(1.0/512.0/0.6, 1.0/512.0/0.6, 1.0/426.0/2.0); - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform); - printTransform(); + glScalef(1.0/S[1]/R[1], 1.0/S[2]/R[2], 1.0/S[3]/R[3]); glTranslatef(position[0], position[1], position[2]); - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform); - printTransform(); glScalef(magnitude[0], magnitude[1], magnitude[0]); - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform); - printTransform(); glRotatef(rot[0], rot[1], rot[2], rot[3]); - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform); - printTransform(); - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform); + float curTrans[16]; + glGetFloatv(GL_TEXTURE_MATRIX, curTrans); + fillTransform(curTrans); + printTransform(); + CHECK_OPENGL_ERROR glMatrixMode(GL_MODELVIEW); } @@ -499,8 +503,8 @@ class gl_spider : public virtual gl_texture setPosition(0.0,0.0,0.0); setDirection(0.0,0.0,1.0); setMagnitude(1.0); - //numSamples = 1089; - numSamples = 9; + numSamples = 1089; + //numSamples = 9; } ///@param samples, the number of samples this spider is going to use. ///best results if samples is can create a perfect root. @@ -538,6 +542,7 @@ class gl_spider : public virtual gl_texture genPositionVectors(); genMagnitudeVectors(); gl_texture::setDims(0.6, 0.6, 2.0); + setSize(512, 512, 426); genTemplate(dirVectors, 0); } @@ -687,15 +692,20 @@ class gl_spider : public virtual gl_texture } void - printTransform() + fillTransform(float mat[16]) { - for(int i = 0; i < 4; i++){ - std::cout << "[" << currentTransform[i] << "]" << - "[" << currentTransform[i+4] << "]" << - "[" << currentTransform[i+8] << "]" << - "[" << currentTransform[i+12] << "]" << - std::endl; + for(int r = 0; r < 4; r++){ + for(int c = 0; c < 4; c++){ + currentTransform(r,c) = mat[c*4+r]; } + } + + } + + void + printTransform() + { + std::cout << currentTransform << std::endl; } /* Method for initializing the cuda devices, necessary only diff --git a/stim/math/vector.h b/stim/math/vector.h index 86fbbfd..2e0c8c7 100644 --- a/stim/math/vector.h +++ b/stim/math/vector.h @@ -203,7 +203,6 @@ struct vec v[i] = rhs; return *this; } - template CUDA_CALLABLE vec & operator=(vec rhs){ for(int i=0; i