From 4166e97375f81aae2bdd5b74c039317b6f1f930a Mon Sep 17 00:00:00 2001 From: pgovyadi Date: Fri, 2 Sep 2016 19:25:05 -0500 Subject: [PATCH] Added debug tags to everything that compile only when the debug option is done, fixed the instability as well as some other issue that came out later --- stim/cuda/testKernel.cuh | 76 ++++++++++++++++++---------------------------------------------------------- stim/gl/gl_spider.h | 123 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++------------------------ 2 files changed, 117 insertions(+), 82 deletions(-) diff --git a/stim/cuda/testKernel.cuh b/stim/cuda/testKernel.cuh index 3e318f8..7ba4ac4 100644 --- a/stim/cuda/testKernel.cuh +++ b/stim/cuda/testKernel.cuh @@ -8,7 +8,7 @@ #include #include #include - stim::cuda::cuda_texture tx; //texture object. + float* print; ///Initialization function, allocates the memory and passes the necessary @@ -26,16 +26,16 @@ cudaFree(print); ///temporary } - __device__ - float templ(int x) - { - if(x < 32/6 || x > 32*5/6 || (x > 32*2/6 && x < 32*4/6)){ - return 1.0; - }else{ - return 0.0; - } - - } + __device__ + float templ(int x, int max_x) + { + if(x < max_x/6 || x > max_x*5/6 || (x > max_x*2/6 && x < max_x*4/6)) + { + return 1.0; + }else{ + return 0.0; + } + } ///Find the difference of the given set of samples and the template ///using cuda acceleration. @@ -44,33 +44,24 @@ ///@param float* result --a pointer to the memory that stores the result. __global__ //void get_diff (float *result) - void get_diff (cudaTextureObject_t texIn, float *print) + void get_diff (cudaTextureObject_t texIn, float *print, int dx) { int x = threadIdx.x + blockIdx.x * blockDim.x; int y = threadIdx.y + blockIdx.y * blockDim.y; -// int idx = y*64+x; - int idx = y*32+x; + int idx = y*dx+x; // int idx = y*16+x; float valIn = tex2D(texIn, x, y); - float templa = templ(x); - print[idx] = valIn; ///temporary + float templa = templ(x, 32)*255.0; + print[idx] = abs(valIn-templa); ///temporary //print[idx] = abs(templa); ///temporary } - - ///External access-point to the cuda function - ///@param GLuint texbufferID --GLtexture (most be contained in a framebuffer object) - /// that holds the data that will be handed to cuda. - ///@param GLenum texType --either GL_TEXTURE_1D, GL_TEXTURE_2D or GL_TEXTURE_3D - /// may work with other gl texture types, but untested. - ///@param DIM_Y, the number of samples in the template. - void test(GLint texbufferID, GLenum texType, int x, int y) + void test(cudaTextureObject_t tObj, int x, int y, std::string nam) { //Bind the Texture in GL and allow access to cuda. - tx.MapCudaTexture(texbufferID, texType); //initialize the return arrays. @@ -85,44 +76,13 @@ // get_diff <<< blocks, threads >>> (tx.getTexture(), print); - get_diff <<< numBlocks, threadsPerBlock >>> (tx.getTexture(), print); - - cudaDeviceSynchronize(); - stringstream name; //for debugging - name << "FromTex.bmp"; - stim::gpu2image(print, name.str(),x,y,0,255); - - tx.UnmapCudaTexture(); - cleanUP(); - } - - void test(GLint texbufferID, GLenum texType, int x, int y, std::string nam) - { - - //Bind the Texture in GL and allow access to cuda. - tx.MapCudaTexture(texbufferID, texType); - - //initialize the return arrays. - - initArray(x,y); - dim3 numBlocks(1, y); - dim3 threadsPerBlock(x, 1); - int max_threads = stim::maxThreadsPerBlock(); - //dim3 threads(max_threads, 1); - //dim3 blocks(x / threads.x + 1, y); - //dim3 numBlocks(2, 2); - //dim3 threadsPerBlock(8, 108); - - -// get_diff <<< blocks, threads >>> (tx.getTexture(), print); - get_diff <<< numBlocks, threadsPerBlock >>> (tx.getTexture(), print); + get_diff <<< numBlocks, threadsPerBlock >>> (tObj, print, x); cudaDeviceSynchronize(); stringstream name; //for debugging name << nam.c_str(); - //stim::gpu2image(print, name.str(),x,y,0,255); + stim::gpu2image(print, name.str(),x,y,0,255); - tx.UnmapCudaTexture(); cleanUP(); } diff --git a/stim/gl/gl_spider.h b/stim/gl/gl_spider.h index f3a007f..73a5f37 100644 --- a/stim/gl/gl_spider.h +++ b/stim/gl/gl_spider.h @@ -134,7 +134,13 @@ class gl_spider // : public virtual gl_texture stim::cuda::cuda_texture t_mag; //cuda_texture object used as an interface between OpenGL and cuda for size vectors. - int iter; + #ifdef DEBUG + stringstream name; + int iter; + int iter_pos; + int iter_dir; + int iter_siz; + #endif //--------------------------------------------------------------------------// //-------------------------------PRIVATE METHODS----------------------------// @@ -155,11 +161,14 @@ class gl_spider // : public virtual gl_texture #ifdef TIMING direction_time += gpuStopTimer(); //profiling #endif - #ifdef TESTING -// test(texbufferID, GL_TEXTURE_2D,2*n_pixels,numSamples*n_pixels, "Final_Cost_Direction.bmp"); - #endif int best = getCost(t_dir.getTexture(), t_dir.getAuxArray() ,numSamples); //find min cost. + #ifdef DEBUG + name.str(""); + name << "Final_Cost_Direction_fiber_"<< iter << "_" << iter_dir << ".bmp"; + test(t_dir.getTexture(), n_pixels*2.0, numSamples*n_pixels, name.str()); + iter_dir++; + #endif stim::vec next( ///calculate the next vector. dV[best][0]*S[0]*R[0], dV[best][1]*S[1]*R[1], @@ -189,10 +198,13 @@ class gl_spider // : public virtual gl_texture position_time += gpuStopTimer(); ///timer for profiling #endif - #ifdef TESTING -// test(texbufferID, GL_TEXTURE_2D,2*n_pixels,numSamples*n_pixels, "Final_Cost_Direction.bmp"); - #endif int best = getCost(t_pos.getTexture(), t_pos.getAuxArray(), numSamplesPos); //find min cost. + #ifdef DEBUG + name.str(""); + name << "Final_Cost_Position_" << iter << "_" << iter_pos << ".bmp"; + test(t_pos.getTexture(), n_pixels*2.0, numSamplesPos*n_pixels, name.str()); + iter_pos++; + #endif stim::vec next( //find next position. pV[best][0], pV[best][1], @@ -221,9 +233,13 @@ class gl_spider // : public virtual gl_texture #ifdef TIMING size_time += gpuStopTimer(); #endif - #ifdef TESTING - #endif int best = getCost(t_mag.getTexture(), t_mag.getAuxArray(), numSamplesMag); //get best cost. + #ifdef DEBUG + name.str(""); + name << "Final_Cost_Size_" << iter << "_" << iter_siz << ".bmp"; + test(t_mag.getTexture(), n_pixels*2.0, numSamplesMag*n_pixels, name.str()); + iter_siz++; + #endif setMagnitude(m*mV[best]); //adjust the magnitude. } @@ -242,7 +258,6 @@ class gl_spider // : public virtual gl_texture if(cL.size() < 4){} ///if the size of the fiber is less then 4 we do nothing. else{ setMatrix(1); ///finds the current transformation matrix - DrawLongCylinder(n, l_template, l_square); ///Draw the cylinder. stim::cylinder cyl(cL, cM); std::vector< stim::vec > result = find_branch(cylinder_texID, GL_TEXTURE_2D, n*l_square, (cL.size()-1)*l_template); ///find all the centers in cuda @@ -269,8 +284,7 @@ class gl_spider // : public virtual gl_texture } stim::vec3 v = cyl.surf(pval, result[i][0]); ///find the coordinates of the point at pval on the surface in tissue space. stim::vec3 di = cyl.p(pval); ///find the coord of v in tissue space projected on the centerline. - float rad = cyl.r(pval); ///find the radius at the pvalue's location - std::cout << v << rad << std::endl; + float rad = cyl.r(pval)/2; ///find the radius at the pvalue's location if( !(v[0] > size[0] || v[1] > size[1] || v[2] > size[2] || v[0] < 0 @@ -355,15 +369,23 @@ class gl_spider // : public virtual gl_texture stim::vec3 sph(1, theta, phi); ///combine into a vector in spherical coordinates. stim::vec3 cart = sph.sph2cart();///convert to cartesian. dV.push_back(cart); ///save the generated vector for further use. + #ifdef DEBUG +// std::cout << cart << std::endl; + #endif if(cos(Y.dot(cart)) < 0.087) ///make sure that the Y is not parallel to the new vector. { Y[0] = 0.0; Y[1] = 1.0; }else{ Y[0] = 1.0; Y[1] = 0.0; } + hor = stim::rect(mag, ///generate a rectangle with the new vectro as a normal. pos, cart, ((Y.cross(cart)).cross(cart)).norm()); + + #ifdef DEBUG + // std::cout << hor.n() << std::endl; + #endif ver = stim::rect(mag, ///generate another rectangle that's perpendicular the first but parallel to the cart vector. pos, cart, hor.n()); @@ -389,10 +411,19 @@ class gl_spider // : public virtual gl_texture //Set up the variable necessary for vector creation. glNewList(dList+1, GL_COMPILE); ///generate a new display list. - for(int i = 0; i < numSamplesPos; i++) ///for the number of position samples + pV.push_back(pos); + hor = stim::rect(mag, ///generate a rec tangle with the new vector as a normal. + pos, dir, + ((Y.cross(d)).cross(d)) + .norm()); + ver = stim::rect(mag, ///generate anoth er rectangle that's perpendicular the first but parallel to the cart vector. + pos, dir, + hor.n()); + UpdateBuffer(0.0, 0.0+0*n_pixels); + for(int i = 1; i < numSamplesPos; i++) ///for the number of position samples { stim::vec3 temp = uniformRandVector(); ///generate a random point on a plane. - temp = temp*delta*2.0 - delta/2.0; ///scale the point + temp = temp*delta*2.0 - delta; ///scale the point and center it around 0.0 temp[2] = 0.0; pV.push_back(temp); ///save the point for further use. hor = stim::rect(mag, ///generate a rectangle with the new vector as a normal. @@ -405,6 +436,10 @@ class gl_spider // : public virtual gl_texture UpdateBuffer(0.0, 0.0+i*n_pixels); ///sample the necessary points and put them into a display list. } glEndList(); ///finilize the display list. + #ifdef DEBUG + for(int i = 0; i < numSamplesPos; i++) + std::cout << pV[i] << std::endl; + #endif } ///@param float delta, How much the rectangles are allowed to expand. @@ -428,7 +463,6 @@ class gl_spider // : public virtual gl_texture float max = 1.0+delta; ///largers size. float step = (max-min)/(numSamplesMag-1); ///the size variation from one rect to the next. float factor; - glNewList(dList+2, GL_COMPILE); for(int i = 0; i < numSamplesMag; i++){ ///for the number of position samples //Create linear index @@ -545,8 +579,8 @@ class gl_spider // : public virtual gl_texture glBindTexture(GL_TEXTURE_2D, textureID); //Textures repeat and use linear interpolation, luminance format. - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_REPEAT); ///Set up the texture to repeat at edges. - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_REPEAT); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); ///Set up the texture to repeat at edges. + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); ///Set up the texture to use Linear interpolation glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR); glTexImage2D(GL_TEXTURE_2D, 0, GL_LUMINANCE, ///Create the texture with no data. @@ -746,7 +780,12 @@ class gl_spider // : public virtual gl_texture numSamples = samples; numSamplesPos = samplespos; numSamplesMag = samplesmag; + #ifdef DEBUG iter = 0; + iter_pos = 0; + iter_dir = 0; + iter_siz = 0; + #endif } ///Position constructor: floats. @@ -770,7 +809,12 @@ class gl_spider // : public virtual gl_texture numSamples = numsamples; numSamplesPos = numsamplespos; numSamplesMag = numsamplesmag; + #ifdef DEBUG iter = 0; + iter_pos = 0; + iter_dir = 0; + iter_siz = 0; + #endif } ///Position constructor: vecs of floats. @@ -789,7 +833,12 @@ class gl_spider // : public virtual gl_texture numSamples = samples; numSamplesPos = samplesPos; numSamplesMag = samplesMag; + #ifdef DEBUG iter = 0; + iter_pos = 0; + iter_dir = 0; + iter_siz = 0; + #endif } ///destructor @@ -825,6 +874,9 @@ class gl_spider // : public virtual gl_texture #endif #ifdef DEBUG iter = 0; + iter_pos = 0; + iter_dir = 0; + iter_siz = 0; #endif stepsize = 3.0; n_pixels = 16.0; @@ -852,7 +904,7 @@ class gl_spider // : public virtual gl_texture genDirectionVectors(5*stim::PI/4); Unbind(); Bind(position_texID, position_buffID, numSamplesPos, n_pixels); - genPositionVectors(); + genPositionVectors(0.2); Unbind(); Bind(radius_texID, radius_buffID, numSamplesMag, n_pixels); genMagnitudeVectors(); @@ -983,17 +1035,20 @@ class gl_spider // : public virtual gl_texture getRotation(stim::vec3 dir) { stim::vec out(0.0,0.0,0.0,0.0); - stim::vec from(0.0,0.0,1.0); + stim::vec3 from(0.0,0.0,1.0); out[0] = acos(dir.dot(from))*180/stim::PI; - if(out[0] < 1.0){ + #ifdef DEBUG + std::cout << "out is " << out << std::endl; + std::cout << "when rotating from " << from << " to " << dir << std::endl; + #endif + if(out[0] < 0.01){ out[0] = 0.0; out[1] = 0.0; out[2] = 0.0; out[3] = 1.0; } else { - stim::vec temp(0.0, 0.0, 0.0);; - stim::vec dir1(dir[0], dir[1], dir[2]); - temp = (from.cross(dir1)).norm(); + stim::vec3 temp(0.0, 0.0, 0.0);; + temp = (from.cross(dir)).norm(); out[1] = temp[0]; out[2] = temp[1]; out[3] = temp[2]; @@ -1226,7 +1281,9 @@ class gl_spider // : public virtual gl_texture int Step() { + #ifdef DEBUG std::cerr << "Took a step" << std::endl; + #endif Bind(direction_texID, direction_buffID, numSamples, n_pixels); CHECK_OPENGL_ERROR findOptimalDirection(); @@ -1376,7 +1433,19 @@ class gl_spider // : public virtual gl_texture setPosition(curSeed); setDirection(curSeedVec); setMagnitude(curSeedMag); - findOptimalScale(); + + #ifdef DEBUG + std::cout << "The new seed " << curSeed << curSeedVec << curSeedMag << std::endl; + #endif + +// Bind(direction_texID, direction_buffID, numSamples, n_pixels); +// CHECK_OPENGL_ERROR +// findOptimalDirection(); +// Unbind(); + +// Bind(radius_texID, radius_buffID, numSamplesMag, n_pixels); +// findOptimalScale(); +// Unbind(); // cL.push_back(curSeed); // cM.push_back(curSeedMag); @@ -1552,6 +1621,9 @@ class gl_spider // : public virtual gl_texture sk.End(); nt.addEdge(L,M); + #ifdef DEBUG + iter++; + #endif } } @@ -1646,6 +1718,9 @@ class gl_spider // : public virtual gl_texture } } } + #ifdef DEBUG + std::cout << "I broke out!" << std::endl; + #endif } }; } -- libgit2 0.21.4