From 2eefb0355e676d2637d3451ae251b1cbcc0df3df Mon Sep 17 00:00:00 2001 From: pgovyadi Date: Wed, 16 Nov 2016 16:54:39 -0600 Subject: [PATCH] added debugging code to the cuda code that is enabled with -DDebug in cmake lists. Fixed minor errors in the grids and image stack classes --- stim/cuda/branch_detection.cuh | 4 ++-- stim/cuda/filter.cuh | 34 +++++++++++++++++++++++++++++----- stim/cuda/testKernel.cuh | 50 +++++++++++++++++++++++++++++++++++++++++++++++++- stim/envi/envi_header.h | 1 + stim/gl/gl_spider.h | 30 ++++++++++++++++++++++-------- stim/gl/gl_texture.h | 19 ++++++++++++------- stim/grids/image_stack.h | 11 ++++++----- 7 files changed, 121 insertions(+), 28 deletions(-) diff --git a/stim/cuda/branch_detection.cuh b/stim/cuda/branch_detection.cuh index 9adf987..33298e3 100644 --- a/stim/cuda/branch_detection.cuh +++ b/stim/cuda/branch_detection.cuh @@ -11,7 +11,7 @@ typedef unsigned int uint; std::vector< stim::vec > -find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y) +find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y, int iter = 0) { float sigma = 2.0; unsigned int conn = 7; @@ -22,7 +22,7 @@ find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y) stringstream name; - cpuCenters = stim::cuda::get_centers(texbufferID, texType, x, y, sizek, sigma, conn, threshold); + cpuCenters = stim::cuda::get_centers(texbufferID, texType, x, y, sizek, sigma, conn, threshold, iter); cudaDeviceSynchronize(); diff --git a/stim/cuda/filter.cuh b/stim/cuda/filter.cuh index 0a0dbf1..497afc7 100644 --- a/stim/cuda/filter.cuh +++ b/stim/cuda/filter.cuh @@ -26,6 +26,11 @@ namespace stim float* LoG; float* res; float* centers; + +//#ifdef DEBUG + float* print; +//#endif + stim::cuda::cuda_texture tx; @@ -44,6 +49,13 @@ namespace stim HANDLE_ERROR( cudaMalloc( (void**) ¢ers, DIM_Y*DIM_X*sizeof(float)) ); + +//#ifdef DEBUG + HANDLE_ERROR( + cudaMalloc( (void**) &print, DIM_Y*DIM_X*sizeof(float)) + ); +//#endif + // checkCUDAerrors("Memory Allocation, Result"); } @@ -58,6 +70,11 @@ namespace stim HANDLE_ERROR( cudaFree(centers) ); +//#ifdef DEBUG + HANDLE_ERROR( + cudaFree(print) + ); +//#endif free(LoG); } @@ -89,7 +106,7 @@ namespace stim //Shared memory would be better. __global__ void - applyFilter(cudaTextureObject_t texIn, unsigned int DIM_X, unsigned int DIM_Y, int kr, int kl, float *res, float* gpuLoG){ + applyFilter(cudaTextureObject_t texIn, unsigned int DIM_X, unsigned int DIM_Y, int kr, int kl, float *res, float* gpuLoG, float* p){ //R = floor(size/2) //THIS IS A NAIVE WAY TO DO IT, and there is a better way) @@ -101,16 +118,15 @@ namespace stim // float val = 0; float tu = (x-kr+xi)/(float)DIM_X; float tv = (y-kr+yi)/(float)DIM_Y; + int idx = y*DIM_X+x; shared[xi][yi] = gpuLoG[yi*kl+xi]*(255.0-(float)tex2D(texIn, tu, tv)); __syncthreads(); - //x = max(0,x); //x = min(x, width-1); //y = max(y, 0); //y = min(y, height - 1); - int idx = y*DIM_X+x; // int k_idx; for(unsigned int step = blockDim.x/2; step >= 1; step >>= 1) { @@ -135,11 +151,12 @@ namespace stim __syncthreads(); if(xi == 0 && yi == 0) res[idx] = shared[0][0]; + } extern "C" float * - get_centers(GLint texbufferID, GLenum texType, int DIM_X, int DIM_Y, int sizeK, float sigma, float conn, float threshold) + get_centers(GLint texbufferID, GLenum texType, int DIM_X, int DIM_Y, int sizeK, float sigma, float conn, float threshold, int iter = 0) { tx.SetTextureCoordinates(1); tx.SetAddressMode(1, 3); @@ -153,7 +170,14 @@ namespace stim dim3 numBlocks(DIM_X, DIM_Y); dim3 threadsPerBlock(sizeK, sizeK); - applyFilter <<< numBlocks, threadsPerBlock >>> (tx.getTexture(), DIM_X, DIM_Y, floor(sizeK/2), sizeK, res, gpuLoG); + applyFilter <<< numBlocks, threadsPerBlock >>> (tx.getTexture(), DIM_X, DIM_Y, floor(sizeK/2), sizeK, res, gpuLoG, print); + + #ifdef DEBUG + stringstream name; + name.str(""); + name << "Fiber Cylinder " << iter << ".bmp"; + stim::gpu2image(res, name.str(), DIM_X, DIM_Y, 0, 255); + #endif stim::cuda::gpu_local_max(centers, res, threshold, conn, DIM_X, DIM_Y); diff --git a/stim/cuda/testKernel.cuh b/stim/cuda/testKernel.cuh index 7ba4ac4..e7a61ad 100644 --- a/stim/cuda/testKernel.cuh +++ b/stim/cuda/testKernel.cuh @@ -53,11 +53,31 @@ float valIn = tex2D(texIn, x, y); float templa = templ(x, 32)*255.0; - print[idx] = abs(valIn-templa); ///temporary + //print[idx] = abs(valIn-templa); ///temporary + print[idx] = abs(valIn); //print[idx] = abs(templa); ///temporary } + ///Find the difference of the given set of samples and the template + ///using cuda acceleration. + ///@param stim::cuda::cuda_texture t --stim texture that holds all the references + /// to the data. + ///@param float* result --a pointer to the memory that stores the result. + __global__ + //void get_diff (float *result) + void get_diff2 (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*dx+x; + // int idx = y*16+x; + + float valIn = tex2D(texIn, x, y); + print[idx] = abs(valIn); ///temporary + + } + void test(cudaTextureObject_t tObj, int x, int y, std::string nam) { @@ -86,3 +106,31 @@ cleanUP(); } + + void test(cudaTextureObject_t tObj, int x, int y, std::string nam, int iter) + { + + //Bind the Texture in GL and allow access to cuda. + + //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_diff2 <<< numBlocks, threadsPerBlock >>> (tObj, print, x); + + cudaDeviceSynchronize(); + stringstream name; //for debugging + name << nam.c_str(); + stim::gpu2image(print, name.str(),x,y,0,255); + + cleanUP(); + } diff --git a/stim/envi/envi_header.h b/stim/envi/envi_header.h index 0ad108c..7669e46 100644 --- a/stim/envi/envi_header.h +++ b/stim/envi/envi_header.h @@ -8,6 +8,7 @@ #include #include #include +#include //information from an ENVI header file //A good resource can be found here: http://www.exelisvis.com/docs/enviheaderfiles.html diff --git a/stim/gl/gl_spider.h b/stim/gl/gl_spider.h index e49edc6..6d3bb80 100644 --- a/stim/gl/gl_spider.h +++ b/stim/gl/gl_spider.h @@ -27,7 +27,6 @@ #include #include "../../../volume-spider/glnetwork.h" #include -#include #include #include #ifdef TIMING @@ -40,6 +39,9 @@ #include #endif +#ifdef DEBUG + #include +#endif namespace stim { @@ -143,8 +145,8 @@ class gl_spider // : public virtual gl_texture #ifdef DEBUG - stringstream name; int iter; + stringstream name; int iter_pos; int iter_dir; int iter_siz; @@ -294,6 +296,7 @@ class gl_spider // : public virtual gl_texture 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 + stim::vec3 size(S[0]*R[0], S[1]*R[1], S[2]*R[2]); ///the borders of the texture. float pval; //pvalue associated with the points on the cylinder. if(!result.empty()) ///if we have any points @@ -957,7 +960,7 @@ class gl_spider // : public virtual gl_texture iter_dir = 0; iter_siz = 0; #endif - stepsize = 3.0; + stepsize = 6.0; n_pixels = 16.0; srand(100); @@ -1380,20 +1383,31 @@ class gl_spider // : public virtual gl_texture Step() { #ifdef DEBUG - std::cerr << "Took a step" << std::endl; + std::cerr << "Took a step"; #endif Bind(direction_texID, direction_buffID, numSamples, n_pixels); CHECK_OPENGL_ERROR findOptimalDirection(); Unbind(); + #ifdef DEBUG + std::cerr << " " << current_cost; + #endif Bind(position_texID, position_buffID, numSamplesPos, n_pixels); findOptimalPosition(); Unbind(); + #ifdef DEBUG + std::cerr << " " << current_cost; + #endif Bind(radius_texID, radius_buffID, numSamplesMag, n_pixels); findOptimalRadius(); Unbind(); + #ifdef DEBUG + std::cerr << " " << current_cost; + #endif CHECK_OPENGL_ERROR - + #ifdef DEBUG + std::cerr << std::endl; + #endif return current_cost; } @@ -1539,9 +1553,9 @@ class gl_spider // : public virtual gl_texture // findOptimalDirection(); // Unbind(); //THIS IS EXPERIMENTAL - Bind(radius_texID, radius_buffID, numSamplesMag, n_pixels); - findOptimalRadius(); - Unbind(); + // Bind(radius_texID, radius_buffID, numSamplesMag, n_pixels); + // findOptimalRadius(); + // Unbind(); //THIS IS EXPERIMENTAL // cL.push_back(curSeed); diff --git a/stim/gl/gl_texture.h b/stim/gl/gl_texture.h index a1314b7..38847c5 100644 --- a/stim/gl/gl_texture.h +++ b/stim/gl/gl_texture.h @@ -30,9 +30,9 @@ class gl_texture : public virtual image_stack GLenum cpu_type; GLenum gpu_type; GLenum format; //format for the texture (GL_RGBA, GL_LUMINANCE, etc.) - using image_stack::R; + using image_stack::R; //using image_stack::S; - using image_stack::ptr; + using image_stack::ptr; /// Sets the internal texture_type, based on the data dimensions void setTextureType(){ @@ -247,7 +247,7 @@ class gl_texture : public virtual image_stack } ///returns the dimentions of the data in the x, y, z directions. - vec getSize(){ + stim::vec getSize(){ stim::vec size(R[1], R[2], R[3]); return size; } @@ -282,7 +282,7 @@ class gl_texture : public virtual image_stack ///@param file_mask specifies the file(s) to be loaded /// Sets the path and calls the loader on that path. void load_images(std::string file_mask){ - image_stack::load_images(file_mask); //load images + image_stack::load_images(file_mask); //load images guess_parameters(); } @@ -292,13 +292,18 @@ class gl_texture : public virtual image_stack return texID; } - + T* getData(){ return ptr; } - - + + void setData(T* rts) + { + + } + }; + } diff --git a/stim/grids/image_stack.h b/stim/grids/image_stack.h index 6bf66a4..0bfef7a 100644 --- a/stim/grids/image_stack.h +++ b/stim/grids/image_stack.h @@ -17,13 +17,13 @@ class image_stack : public virtual stim::grid{ protected: //using stim::grid::S; - using stim::grid::R; - using stim::grid::ptr; - using stim::grid::read; + using stim::grid::R; + using stim::grid::ptr; + using stim::grid::read; public: //default constructor - image_stack() : grid() { + image_stack() : grid() { } @@ -113,7 +113,8 @@ public: /// @param i is the page to be saved void save_image(std::string file_name, unsigned int i){ stim::image I; //create an image - I.set_interleaved_rgb(&ptr[ i * R[0] * R[1] * R[2] ], R[1], R[2], R[0]); //retrieve the interlaced data from the image - store it in the grid + I.set_interleaved(&ptr[ i * R[0] * R[1] * R[2] ], R[1], R[2], R[0]); //retrieve the interlaced data from the image - store it in the grid +// I.set_interleaved_rgb(&ptr[ i * R[0] * R[1] * R[2] ], R[1], R[2], R[0]); //retrieve the interlaced data from the image - store it in the grid I.save(file_name); } -- libgit2 0.21.4