Commit 59071526e1543936be7509ca1d609b986ce0458a

Authored by David Mayerich
2 parents ee52e24f d9a27f16

Merge branch 'master' of git.stim.ee.uh.edu:codebase/stimlib

Showing 2 changed files with 26 additions and 48 deletions   Show diff stats
stim/iVote/ivote2.cuh
@@ -94,8 +94,7 @@ namespace stim { @@ -94,8 +94,7 @@ namespace stim {
94 94
95 //this function performs the 2D iterative voting algorithm on the image stored in the gpu 95 //this function performs the 2D iterative voting algorithm on the image stored in the gpu
96 template<typename T> 96 template<typename T>
97 - void gpu_ivote2(T* gpuI, unsigned int rmax, size_t x, size_t y, bool invert = false, T t = 0, std::string outname_img = "out.bmp", std::string outname_txt = "out.txt",  
98 - int iter = 8, T phi = 15.0f * (float)stim::PI / 180, int conn = 8, bool debug = false) { 97 + void gpu_ivote2(T* gpuI, unsigned int rmax, size_t x, size_t y, bool invert = false, T t = 0, int iter = 8, T phi = 15.0f * (float)stim::PI / 180, int conn = 8, bool debug = false) {
99 98
100 size_t pixels = x * y; //compute the size of input image 99 size_t pixels = x * y; //compute the size of input image
101 // 100 //
@@ -129,43 +128,36 @@ namespace stim { @@ -129,43 +128,36 @@ namespace stim {
129 } 128 }
130 stim::cuda::gpu_local_max<float>(gpuI, gpuVote, conn, x, y); //calculate the local maxima 129 stim::cuda::gpu_local_max<float>(gpuI, gpuVote, conn, x, y); //calculate the local maxima
131 130
132 - T* pts = (T*)malloc(bytes); //allocate memory on the cpu to store the output of iterative voting  
133 - HANDLE_ERROR(cudaMemcpy(pts, gpuI, bytes, cudaMemcpyDeviceToHost)); //copy the output from gpu to the cpu memory  
134 -  
135 - T threshold;  
136 - if (t == 0) threshold = stim::th_otsu<T>(pts, pixels); //if threshold value is not set call the function to compute the threshold  
137 - else threshold = t;  
138 -  
139 - std::ofstream output; //save the thresholded detected seeds in a text file  
140 - output.open(outname_txt);  
141 - output << "X" << " " << "Y" << " " << "threshold" << "\n";  
142 - size_t ind;  
143 - for (size_t ix = 0; ix < x; ix++) {  
144 - for (size_t iy = 0; iy < y; iy++) {  
145 - ind = iy * x + ix;  
146 - if (pts[ind] > threshold) {  
147 - output << ix << " " << iy << " " << pts[ind] << "\n";  
148 - pts[ind] = 1; 131 + if (t > 0) {
  132 + T* pts = (T*)malloc(bytes); //allocate memory on the cpu to store the output of iterative voting
  133 + HANDLE_ERROR(cudaMemcpy(pts, gpuI, bytes, cudaMemcpyDeviceToHost)); //copy the output from gpu to the cpu memory
  134 +
  135 + T threshold;
  136 + threshold = t;
  137 +
  138 + size_t ind;
  139 + for (size_t ix = 0; ix < x; ix++) {
  140 + for (size_t iy = 0; iy < y; iy++) {
  141 + ind = iy * x + ix;
  142 + if (pts[ind] > threshold) {
  143 + pts[ind] = 1;
  144 + }
  145 + else pts[ind] = 0;
149 } 146 }
150 - else pts[ind] = 0;  
151 } 147 }
  148 + HANDLE_ERROR(cudaMemcpy(gpuI, pts, bytes, cudaMemcpyHostToDevice)); //copy the points to the gpu
152 } 149 }
153 - output.close();  
154 -  
155 - HANDLE_ERROR(cudaMemcpy(gpuI, pts, bytes, cudaMemcpyHostToDevice)); //copy the points to the gpu  
156 - stim::cpu2image(pts, outname_img, x, y); //output the image  
157 - 150 +
158 } 151 }
159 152
160 153
161 template<typename T> 154 template<typename T>
162 - void cpu_ivote2(T* cpuI, unsigned int rmax, size_t x, size_t y, bool invert = false, T t = 0, std::string outname_img = "out.bmp", std::string outname_txt = "out.txt",  
163 - int iter = 8, T phi = 15.0f * (float)stim::PI / 180, int conn = 8, bool debug = false) { 155 + void cpu_ivote2(T* cpuI, unsigned int rmax, size_t x, size_t y, bool invert = false, T t = 0, int iter = 8, T phi = 15.0f * (float)stim::PI / 180, int conn = 8, bool debug = false) {
164 size_t bytes = x*y * sizeof(T); 156 size_t bytes = x*y * sizeof(T);
165 T* gpuI; //allocate space on the gpu to save the input image 157 T* gpuI; //allocate space on the gpu to save the input image
166 HANDLE_ERROR(cudaMalloc(&gpuI, bytes)); 158 HANDLE_ERROR(cudaMalloc(&gpuI, bytes));
167 HANDLE_ERROR(cudaMemcpy(gpuI, cpuI, bytes, cudaMemcpyHostToDevice)); //copy the image to the gpu 159 HANDLE_ERROR(cudaMemcpy(gpuI, cpuI, bytes, cudaMemcpyHostToDevice)); //copy the image to the gpu
168 - stim::gpu_ivote2<T>(gpuI, rmax, x, y, invert, t, outname_img, outname_txt, iter, phi, conn, debug); //call the gpu version of the ivote 160 + stim::gpu_ivote2<T>(gpuI, rmax, x, y, invert, t, iter, phi, conn, debug); //call the gpu version of the ivote
169 HANDLE_ERROR(cudaMemcpy(cpuI, gpuI, bytes, cudaMemcpyDeviceToHost)); //copy the output to the cpu 161 HANDLE_ERROR(cudaMemcpy(cpuI, gpuI, bytes, cudaMemcpyDeviceToHost)); //copy the output to the cpu
170 } 162 }
171 } 163 }
stim/iVote/ivote2/local_max.cuh
@@ -10,30 +10,22 @@ namespace stim{ @@ -10,30 +10,22 @@ namespace stim{
10 10
11 // this kernel calculates the local maximum for finding the cell centers 11 // this kernel calculates the local maximum for finding the cell centers
12 template<typename T> 12 template<typename T>
13 - __global__ void cuda_local_max(T* gpuCenters, T* gpuVote, int conn, size_t x, size_t y){ 13 + __global__ void cuda_local_max(T* gpuCenters, T* gpuVote, int conn, int x, int y){
14 14
15 // calculate the 2D coordinates for this current thread. 15 // calculate the 2D coordinates for this current thread.
16 - size_t xi = blockIdx.x * blockDim.x + threadIdx.x;  
17 - size_t yi = blockIdx.y * blockDim.y + threadIdx.y; 16 + int xi = blockIdx.x * blockDim.x + threadIdx.x;
  17 + int yi = blockIdx.y * blockDim.y + threadIdx.y;
18 18
19 if(xi >= x || yi >= y) 19 if(xi >= x || yi >= y)
20 return; 20 return;
21 21
22 // convert 2D coordinates to 1D 22 // convert 2D coordinates to 1D
23 - size_t i = yi * x + xi; 23 + int i = yi * x + xi;
24 24
25 gpuCenters[i] = 0; //initialize the value at this location to zero 25 gpuCenters[i] = 0; //initialize the value at this location to zero
26 26
27 T val = gpuVote[i]; 27 T val = gpuVote[i];
28 28
29 - //compare to the threshold  
30 - //if(val < final_t) return;  
31 -  
32 - //define an array to store indices with same vote value  
33 - /*int * IdxEq;  
34 - IdxEq = new int [2*conn];  
35 - int n = 0;*/  
36 -  
37 for(int xl = xi - conn; xl < xi + conn; xl++){ 29 for(int xl = xi - conn; xl < xi + conn; xl++){
38 for(int yl = yi - conn; yl < yi + conn; yl++){ 30 for(int yl = yi - conn; yl < yi + conn; yl++){
39 if(xl >= 0 && xl < x && yl >= 0 && yl < y){ 31 if(xl >= 0 && xl < x && yl >= 0 && yl < y){
@@ -42,8 +34,7 @@ namespace stim{ @@ -42,8 +34,7 @@ namespace stim{
42 return; 34 return;
43 } 35 }
44 if (gpuVote[il] == val){ 36 if (gpuVote[il] == val){
45 - /*IdxEq[n] = il;  
46 - n = n+1;*/ 37 +
47 if( il > i){ 38 if( il > i){
48 return; 39 return;
49 } 40 }
@@ -51,12 +42,7 @@ namespace stim{ @@ -51,12 +42,7 @@ namespace stim{
51 } 42 }
52 } 43 }
53 } 44 }
54 - /*if (n!=0){  
55 - if(IdxEq[n/2] !=i){  
56 - return;  
57 - }  
58 - } */  
59 - //gpuCenters[i] = 1; 45 +
60 gpuCenters[i] = gpuVote[i]; 46 gpuCenters[i] = gpuVote[i];
61 } 47 }
62 48