Blame view

stim/cuda/ivote/update_dir_global.cuh 9.26 KB
11cd127f   Laila Saadatifard   Leila's ivote pro...
1
2
3
4
5
6
7
  #ifndef STIM_CUDA_UPDATE_DIR_GLOBALD_H
  #define STIM_CUDA_UPDATE_DIR_GLOBAL_H
  
  # include <iostream>
  # include <cuda.h>
  #include <stim/cuda/cudatools.h>
  #include <stim/cuda/sharedmem.cuh>
ca99f951   David Mayerich   faster implementa...
8
9
10
  #include <stim/visualization/aabb2.h>
  #include <stim/visualization/colormap.h>
  #include <math.h>
11cd127f   Laila Saadatifard   Leila's ivote pro...
11
12
  #include "cpyToshare.cuh" 
  
ca99f951   David Mayerich   faster implementa...
13
  //#define RMAX_TEST	8
11cd127f   Laila Saadatifard   Leila's ivote pro...
14
15
16
  
  namespace stim{
  	namespace cuda{
ca99f951   David Mayerich   faster implementa...
17
  
11cd127f   Laila Saadatifard   Leila's ivote pro...
18
19
  		template<typename T>
  		__global__ void cuda_update_dir(T* gpuDir, T* gpuVote, T* gpuGrad, T* gpuTable, T phi, int rmax,  int x,  int y){
ca99f951   David Mayerich   faster implementa...
20
21
22
23
24
25
26
27
28
29
30
31
32
33
  			extern __shared__ T S[];
  			T* shared_atan = S;
  			size_t n_table = (rmax * 2 + 1) * (rmax * 2 + 1);
  			stim::cuda::threadedMemcpy((char*)shared_atan, (char*)gpuTable, sizeof(T) * n_table, threadIdx.x, blockDim.x);
  
  			//T* shared_vote = &S[n_table];
  			//size_t template_size_x = (blockDim.x + 2 * rmax);
  			//size_t template_size_y = (blockDim.y + 2 * rmax);
  			//stim::cuda::threadedMemcpy2D((char*)shared_vote, (char*)gpuVote, template_size_x, template_size_y, x,  threadIdx.y * blockDim.x + threadIdx.x, blockDim.x * blockDim.y);
  			
  			int xi = blockIdx.x * blockDim.x + threadIdx.x;				//calculate the 2D coordinates for this current thread.
  			int yi = blockIdx.y * blockDim.y + threadIdx.y;
  
  			if(xi >= x || yi >= y) return;								//if the index is outside of the image, terminate the kernel
11cd127f   Laila Saadatifard   Leila's ivote pro...
34
  
ca99f951   David Mayerich   faster implementa...
35
36
  			int i = yi * x + xi;										//convert 2D coordinates to 1D
  			float theta = gpuGrad[2*i];									//calculate the voting direction based on the grtadient direction - global memory fetch
11cd127f   Laila Saadatifard   Leila's ivote pro...
37
  			
ca99f951   David Mayerich   faster implementa...
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
  			stim::aabb2<int> bb(xi, yi);								//initialize a bounding box at the current point
  			bb.insert(xi + ceil(rmax * cos(theta)),       ceil(yi + rmax * sin(theta)));
  			bb.insert(xi + ceil(rmax * cos(theta - phi)), yi + ceil(rmax * sin(theta - phi)));		//insert one corner of the triangle into the bounding box
  			bb.insert(xi + ceil(rmax * cos(theta + phi)), yi + ceil(rmax * sin(theta + phi)));		//insert the final corner into the bounding box
  
  			int x_table = 2*rmax +1;
  			int lut_i;
  			T rmax_sq = rmax * rmax;
  			T dx_sq, dy_sq;
  
  			bb.trim_low(0, 0);															//make sure the bounding box doesn't go outside the image
  			bb.trim_high(x-1, y-1);
  
  			int by, bx;
  			int dx, dy;													//coordinate relative to (xi, yi)
  			T v;
  			T max_v = 0;												//initialize the maximum vote value to zero
  			T alpha;
  			int max_dx = bb.low[0];
  			int max_dy = bb.low[1];
  			for(by = bb.low[1]; by <= bb.high[1]; by++){					//for each element in the bounding box
  				dy = by - yi;											//calculate the y coordinate of the current point relative to yi
  				dy_sq = dy * dy;
  				for(bx = bb.low[0]; bx <= bb.high[0]; bx++){
  					dx = bx - xi;
  					dx_sq = dx * dx;
  					lut_i = (rmax - dy) * x_table + rmax - dx;
  					alpha = shared_atan[lut_i];
  					if(dx_sq + dy_sq < rmax_sq && abs(alpha - theta) < phi){
  						v = gpuVote[by * x + bx];				// find the vote value for the current counter
  						if(v > max_v){
  							max_v = v;
  							max_dx = dx;
  							max_dy = dy;
  						}
  					}
  				}
  			}			
  			gpuDir[i] = atan2((T)max_dy, (T)max_dx);
  		}
  	
  		// this kernel calculates the voting direction for the next iteration based on the angle between the location of this voter and the maximum vote value in its voting area.
  		template<typename T>
  		__global__ void leila_cuda_update_dir(T* gpuDir, T* gpuVote, T* gpuGrad, T* gpuTable, T phi, int rmax,  int x,  int y){
  
11cd127f   Laila Saadatifard   Leila's ivote pro...
83
84
  			
  			// calculate the 2D coordinates for this current thread.
ca99f951   David Mayerich   faster implementa...
85
86
87
88
89
  			int xi = blockIdx.x * blockDim.x + threadIdx.x;
  			int yi = blockIdx.y * blockDim.y + threadIdx.y;
  
  			if(xi >= x || yi >= y) return;													//if the index is outside of the image, terminate the kernel
  
11cd127f   Laila Saadatifard   Leila's ivote pro...
90
91
92
93
94
95
96
97
98
99
100
101
102
  			int i = yi * x + xi;												// convert 2D coordinates to 1D
  			
  			float theta = gpuGrad[2*i];											// calculate the voting direction based on the grtadient direction - global memory fetch			
  			gpuDir[i] = 0;														//initialize the vote direction to zero			
  			float max = 0;														// define a local variable to maximum value of the vote image in the voting area for this voter
  			int id_x = 0;														// define two local variables for the x and y position of the maximum
  			int id_y = 0;
  			
  			int x_table = 2*rmax +1;											// compute the size of window which will be checked for finding the voting area for this voter
  			int rmax_sq = rmax * rmax;
  			int tx_rmax = threadIdx.x + rmax;
  			float atan_angle;
  			float vote_c;
ca99f951   David Mayerich   faster implementa...
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
  			int xidx, yidx, yr_sq, xr_sq;
  			for(int yr = -rmax; yr <= rmax; yr++){
  				yidx = yi + yr;													//compute the index into the image
  				if (yidx >= 0 && yidx < y){									//if the current y-index is inside the image
  					yr_sq = yr * yr;											//compute the square of yr, to save time later
  					for(int xr = -rmax; xr <= rmax; xr++){
  						xidx = xi + xr;
  						if(xidx >= 0 && xidx < x){
  							xr_sq = xr * xr;
  							unsigned int ind_t = (rmax - yr) * x_table + rmax - xr;
  
  							// calculate the angle between the voter and the current pixel in x and y directions
  							atan_angle = gpuTable[ind_t];
  							//atan_angle = atan2((T)yr, (T)xr);
  											
  							// check if the current pixel is located in the voting area of this voter.
  							if (((xr_sq + yr_sq)< rmax_sq) && (abs(atan_angle - theta) <phi)){
  								
  								vote_c = gpuVote[yidx * x + xidx];				// find the vote value for the current counter
  							// compare the vote value of this pixel with the max value to find the maxima and its index.
  								if  (vote_c>max) {
  
  									max = vote_c;
  									id_x =  xr;
  									id_y =  yr;
  								}
11cd127f   Laila Saadatifard   Leila's ivote pro...
129
130
131
132
133
134
135
136
137
138
139
140
  							}
  						}
  					}
  				}
  			}
  							
  			unsigned int ind_m = (rmax - id_y) * x_table + (rmax - id_x);
  			float new_angle = gpuTable[ind_m];
  
  			if(xi < x && yi < y)
  				gpuDir[i] = new_angle;
  		}										//end kernel
ca99f951   David Mayerich   faster implementa...
141
  		
11cd127f   Laila Saadatifard   Leila's ivote pro...
142
143
144
145
146
147
148
149
  
  		// this kernel updates the gradient direction by the calculated voting direction.
  		template<typename T>
  		__global__ void cuda_update_grad(T* gpuGrad, T* gpuDir, int x, int y){
  
  			// calculate the 2D coordinates for this current thread.
  			int xi = blockIdx.x * blockDim.x + threadIdx.x;
  			int yi = blockIdx.y * blockDim.y + threadIdx.y;
ca99f951   David Mayerich   faster implementa...
150
151
  
  			if(xi >= x || yi >= y) return;
11cd127f   Laila Saadatifard   Leila's ivote pro...
152
153
154
155
156
157
158
159
160
161
162
163
164
  		
  			// convert 2D coordinates to 1D
  			int i = yi * x + xi;
  			
  			//update the gradient image with the vote direction
  			gpuGrad[2*i] = gpuDir[i];
  		}
  		
  		template<typename T>
  		void gpu_update_dir(T* gpuVote, T* gpuGrad, T* gpuTable, T phi, unsigned int rmax, unsigned int x, unsigned int y){
  
  			//calculate the number of bytes in the array
  			unsigned int bytes = x * y * sizeof(T);
11cd127f   Laila Saadatifard   Leila's ivote pro...
165
166
167
  			
  			// allocate space on the GPU for the updated vote direction
  			T* gpuDir;
ca99f951   David Mayerich   faster implementa...
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
  			HANDLE_ERROR( cudaMalloc(&gpuDir, bytes) );	
  
  			unsigned int max_threads = stim::maxThreadsPerBlock();
  			//dim3 threads(min(x, max_threads), 1);
  			//dim3 blocks(x/threads.x, y);
  
  			dim3 threads( sqrt(max_threads), sqrt(max_threads) );
  			dim3 blocks(x/threads.x + 1, y/threads.y + 1);
  
  			size_t table_bytes = sizeof(T) * (rmax * 2 + 1) * (rmax * 2 + 1);
  			//size_t curtain = 2 * rmax;
  			//size_t template_bytes = sizeof(T) * (threads.x + curtain) * (threads.y + curtain);
  			size_t shared_mem_req = table_bytes;// + template_bytes;
  			std::cout<<"Shared Memory required: "<<shared_mem_req<<std::endl;
  
  			size_t shared_mem = stim::sharedMemPerBlock();
  			if(shared_mem_req > shared_mem){
  				std::cout<<"Error: insufficient shared memory for this implementation of cuda_update_dir()."<<std::endl;
  				exit(1);
  			}
11cd127f   Laila Saadatifard   Leila's ivote pro...
188
189
  
  			//call the kernel to calculate the new voting direction
ca99f951   David Mayerich   faster implementa...
190
191
192
193
194
195
196
  			cuda_update_dir <<< blocks, threads, shared_mem_req>>>(gpuDir, gpuVote, gpuGrad, gpuTable, phi, rmax, x , y);
  			stim::gpu2image<T>(gpuDir, "dir_david.bmp", x, y, -pi, pi, stim::cmBrewer);
  
  			//exit(0);
  
  			threads = dim3( sqrt(max_threads), sqrt(max_threads) );
  			blocks = dim3(x/threads.x + 1, y/threads.y + 1);
11cd127f   Laila Saadatifard   Leila's ivote pro...
197
198
199
  
  			//call the kernel to update the gradient direction
  			cuda_update_grad <<< blocks, threads >>>(gpuGrad, gpuDir, x , y);
11cd127f   Laila Saadatifard   Leila's ivote pro...
200
  			//free allocated memory
ca99f951   David Mayerich   faster implementa...
201
  			HANDLE_ERROR( cudaFree(gpuDir) );
11cd127f   Laila Saadatifard   Leila's ivote pro...
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
  
  		}
  		
  		template<typename T>
  		void cpu_update_dir(T* cpuVote, T* cpuGrad,T* cpuTable, T phi, unsigned int rmax, unsigned int x, unsigned int y){
  
  			//calculate the number of bytes in the array
  			unsigned int bytes = x * y * sizeof(T);
  
  			//calculate the number of bytes in the atan2 table
  			unsigned int bytes_table = (2*rmax+1) * (2*rmax+1) * sizeof(T);
  
  			//allocate space on the GPU for the Vote Image
  			T* gpuVote;
  			cudaMalloc(&gpuVote, bytes);
  
  			//copy the input vote image to the GPU
  			HANDLE_ERROR(cudaMemcpy(gpuVote, cpuVote, bytes, cudaMemcpyHostToDevice));	
  
  			//allocate space on the GPU for the input Gradient image
  			T* gpuGrad;
  			HANDLE_ERROR(cudaMalloc(&gpuGrad, bytes*2));
  
  			//copy the Gradient data to the GPU
  			HANDLE_ERROR(cudaMemcpy(gpuGrad, cpuGrad, bytes*2, cudaMemcpyHostToDevice));
  
  			//allocate space on the GPU for the atan2 table
  			T* gpuTable;
  			HANDLE_ERROR(cudaMalloc(&gpuTable, bytes_table));
  
  			//copy the atan2 values to the GPU
  			HANDLE_ERROR(cudaMemcpy(gpuTable, cpuTable, bytes_table, cudaMemcpyHostToDevice));
  						
  			//call the GPU version of the update direction function
  			gpu_update_dir<T>(gpuVote, gpuGrad, gpuTable, phi, rmax, x , y);
  							
  			//copy the new gradient image back to the CPU
  			cudaMemcpy(cpuGrad, gpuGrad, bytes*2, cudaMemcpyDeviceToHost) ;
  
  			//free allocated memory
  			cudaFree(gpuTable);
  			cudaFree(gpuVote);
  			cudaFree(gpuGrad);
  		}
  		
  	}
  }
  
  #endif