Blame view

stim/iVote/ivote2/update_dir_bb.cuh 6.41 KB
3f0de7dd   Laila Saadatifard   upload the vote a...
1
2
  #ifndef STIM_CUDA_UPDATE_DIR_BB_H
  #define STIM_CUDA_UPDATE_DIR_BB_H
11cd127f   Laila Saadatifard   Leila's ivote pro...
3
4
5
6
7
  
  # include <iostream>
  # include <cuda.h>
  #include <stim/cuda/cudatools.h>
  #include <stim/cuda/sharedmem.cuh>
ca99f951   David Mayerich   faster implementa...
8
9
  #include <stim/visualization/aabb2.h>
  #include <stim/visualization/colormap.h>
03428452   Laila Saadatifard   update the ivote ...
10
  #include <math.h> 
11cd127f   Laila Saadatifard   Leila's ivote pro...
11
  
ca99f951   David Mayerich   faster implementa...
12
  //#define RMAX_TEST	8
11cd127f   Laila Saadatifard   Leila's ivote pro...
13
14
15
  
  namespace stim{
  	namespace cuda{
ca99f951   David Mayerich   faster implementa...
16
  
11cd127f   Laila Saadatifard   Leila's ivote pro...
17
18
  		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...
19
20
21
22
23
24
25
26
27
28
29
30
31
32
  			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...
33
  
ca99f951   David Mayerich   faster implementa...
34
35
  			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...
36
  			
ca99f951   David Mayerich   faster implementa...
37
38
39
40
41
42
  			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;
ca99f951   David Mayerich   faster implementa...
43
  			T rmax_sq = rmax * rmax;
daacc99c   Laila Saadatifard   change the local ...
44
45
  
  			int lut_i;
ca99f951   David Mayerich   faster implementa...
46
47
48
49
50
51
52
  			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)
daacc99c   Laila Saadatifard   change the local ...
53
  			
ca99f951   David Mayerich   faster implementa...
54
55
56
  			T v;
  			T max_v = 0;												//initialize the maximum vote value to zero
  			T alpha;
daacc99c   Laila Saadatifard   change the local ...
57
58
  			int max_dx = bb.low[0] - xi;
  			int max_dy = bb.low[1] - yi;
ca99f951   David Mayerich   faster implementa...
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
  			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);
  		}
  	
ca99f951   David Mayerich   faster implementa...
80
  		
11cd127f   Laila Saadatifard   Leila's ivote pro...
81
82
83
  
  		// this kernel updates the gradient direction by the calculated voting direction.
  		template<typename T>
276f9b23   David Mayerich   forced matching t...
84
  		__global__ void cuda_update_grad(T* gpuGrad, T* gpuDir, size_t x, size_t y){
11cd127f   Laila Saadatifard   Leila's ivote pro...
85
86
  
  			// calculate the 2D coordinates for this current thread.
276f9b23   David Mayerich   forced matching t...
87
88
  			size_t xi = blockIdx.x * blockDim.x + threadIdx.x;
  			size_t yi = blockIdx.y * blockDim.y + threadIdx.y;
ca99f951   David Mayerich   faster implementa...
89
90
  
  			if(xi >= x || yi >= y) return;
11cd127f   Laila Saadatifard   Leila's ivote pro...
91
92
  		
  			// convert 2D coordinates to 1D
276f9b23   David Mayerich   forced matching t...
93
  			size_t i = yi * x + xi;
11cd127f   Laila Saadatifard   Leila's ivote pro...
94
95
96
97
98
99
  			
  			//update the gradient image with the vote direction
  			gpuGrad[2*i] = gpuDir[i];
  		}
  		
  		template<typename T>
6316fca0   Laila Saadatifard   fix the 'unresolv...
100
  		void gpu_update_dir(T* gpuVote, T* gpuGrad, T* gpuTable, T phi, unsigned int rmax, size_t x, size_t y, bool DEBUG = false){
11cd127f   Laila Saadatifard   Leila's ivote pro...
101
102
  
  			//calculate the number of bytes in the array
276f9b23   David Mayerich   forced matching t...
103
  			size_t bytes = x * y * sizeof(T);
11cd127f   Laila Saadatifard   Leila's ivote pro...
104
105
106
  			
  			// allocate space on the GPU for the updated vote direction
  			T* gpuDir;
ca99f951   David Mayerich   faster implementa...
107
108
109
  			HANDLE_ERROR( cudaMalloc(&gpuDir, bytes) );	
  
  			unsigned int max_threads = stim::maxThreadsPerBlock();
03428452   Laila Saadatifard   update the ivote ...
110
  			
276f9b23   David Mayerich   forced matching t...
111
112
  			dim3 threads( (unsigned int)sqrt(max_threads), (unsigned int)sqrt(max_threads) );
  			dim3 blocks((unsigned int)x/threads.x + 1, (unsigned int)y/threads.y + 1);
ca99f951   David Mayerich   faster implementa...
113
114
115
116
117
  
  			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;
276f9b23   David Mayerich   forced matching t...
118
  			if (DEBUG) std::cout << "Shared Memory required: " << shared_mem_req << std::endl;
ca99f951   David Mayerich   faster implementa...
119
120
121
122
123
124
  
  			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...
125
126
  
  			//call the kernel to calculate the new voting direction
276f9b23   David Mayerich   forced matching t...
127
  			cuda_update_dir <<< blocks, threads, shared_mem_req>>>(gpuDir, gpuVote, gpuGrad, gpuTable, phi, rmax, (int)x , (int)y);
11cd127f   Laila Saadatifard   Leila's ivote pro...
128
129
  
  			//call the kernel to update the gradient direction
276f9b23   David Mayerich   forced matching t...
130
  			cuda_update_grad <<< blocks, threads >>>(gpuGrad, gpuDir, (int)x , (int)y);
11cd127f   Laila Saadatifard   Leila's ivote pro...
131
  			//free allocated memory
ca99f951   David Mayerich   faster implementa...
132
  			HANDLE_ERROR( cudaFree(gpuDir) );
11cd127f   Laila Saadatifard   Leila's ivote pro...
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
  
  		}
  		
  		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