13fe3c84
 Laila Saadatifard
update the stimli...
|
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
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
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
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
129
130
131
132
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
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
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
|
#include <stim/cuda/sharedmem.cuh>
namespace stim{
namespace cuda{
// 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 cuda_update_dir(T* gpuDir, cudaTextureObject_t in, T* gpuGrad, T* gpuTable, T phi, unsigned int rmax, unsigned int x, unsigned int y){
//generate a pointer to shared memory (size will be specified as a kernel parameter)
extern __shared__ float s_vote[];
//calculate the start point for this block
int bxi = blockIdx.x * blockDim.x;
//calculate the width of the shared memory block
int swidth = 2 * rmax + blockDim.x;
// calculate the 2D coordinates for this current thread.
int xi = bxi + threadIdx.x;
int yi = blockIdx.y;
// convert 2D coordinates to 1D
int i = yi * x + xi;
// calculate the voting direction based on the grtadient direction
float theta = gpuGrad[2*i];
//initialize the vote direction to zero
gpuDir[i] = 0;
// define a local variable to maximum value of the vote image in the voting area for this voter
float max = 0;
// define two local variables for the x and y coordinations where the maximum happened
int id_x = 0;
int id_y = 0;
// compute the size of window which will be checked for finding the voting area for this voter
unsigned int x_table = 2*rmax +1;
unsigned int rmax_sq = rmax * rmax;
int r = (int)rmax;
int tx_rmax = threadIdx.x + rmax;
int bxs = bxi - rmax;
for(int yr = -r; yr <= r; yr++){
//copy the portion of the image necessary for this block to shared memory
__syncthreads();
stim::cuda::sharedMemcpy_tex2D<float>(s_vote, in, bxs, yi + yr , swidth, 1, threadIdx, blockDim);
__syncthreads();
//if the current thread is outside of the image, it doesn't have to be computed
if(xi < x && yi < y){
for(int xr = -r; xr <= r; 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
float atan_angle = gpuTable[ind_t];
// calculate the voting direction based on the grtadient direction
int idx_share_update = xr + tx_rmax ;
float share_vote = s_vote[idx_share_update];
// check if the current pixel is located in the voting area of this voter.
if (((xr * xr + yr *yr)< rmax_sq) && (abs(atan_angle - theta) <phi)){
// compare the vote value of this pixel with the max value to find the maxima and its index.
if (share_vote>max) {
max = share_vote;
id_x = xr;
id_y = yr;
}
}
}
}
}
//float new_angle = atan2(dy, dx);
unsigned int ind_m = (rmax - id_y) * x_table + (rmax - id_x);
float new_angle = gpuTable[ind_m];
gpuDir[i] = new_angle;
}
// this kernel updates the gradient direction by the calculated voting direction.
template<typename T>
__global__ void cuda_update_grad(T* gpuGrad, T* gpuDir, unsigned int x, unsigned int y){
//************ when the number of threads are (1024,1) *************
// calculate the 2D coordinates for this current thread.
int xi = blockIdx.x * blockDim.x + threadIdx.x;
int yi = blockIdx.y;
// 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){
//get the number of pixels in the image
unsigned int pixels = x * y;
unsigned int bytes = sizeof(T) * pixels;
unsigned int max_threads = stim::maxThreadsPerBlock();
dim3 threads(max_threads, 1);
dim3 blocks(x/threads.x + (x %threads.x == 0 ? 0:1) , y);
// Allocate CUDA array in device memory
//define a channel descriptor for a single 32-bit channel
cudaChannelFormatDesc channelDesc =
cudaCreateChannelDesc(32, 0, 0, 0,
cudaChannelFormatKindFloat);
cudaArray* cuArray; //declare the cuda array
cudaMallocArray(&cuArray, &channelDesc, x, y); //allocate the cuda array
// Copy the image data from global memory to the array
cudaMemcpyToArray(cuArray, 0, 0, gpuVote, bytes,
cudaMemcpyDeviceToDevice);
// Specify texture
struct cudaResourceDesc resDesc; //create a resource descriptor
memset(&resDesc, 0, sizeof(resDesc)); //set all values to zero
resDesc.resType = cudaResourceTypeArray; //specify the resource descriptor type
resDesc.res.array.array = cuArray; //add a pointer to the cuda array
// Specify texture object parameters
struct cudaTextureDesc texDesc; //create a texture descriptor
memset(&texDesc, 0, sizeof(texDesc)); //set all values in the texture descriptor to zero
texDesc.addressMode[0] = cudaAddressModeWrap; //use wrapping (around the edges)
texDesc.addressMode[1] = cudaAddressModeWrap;
texDesc.filterMode = cudaFilterModePoint; //use linear filtering
texDesc.readMode = cudaReadModeElementType; //reads data based on the element type (32-bit floats)
texDesc.normalizedCoords = 0; //not using normalized coordinates
// Create texture object
cudaTextureObject_t texObj = 0;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);
// specify share memory
unsigned int share_bytes = (2*rmax + threads.x)*(1)*4;
// allocate space on the GPU for the updated vote direction
T* gpuDir;
cudaMalloc(&gpuDir, bytes);
//call the kernel to calculate the new voting direction
cuda_update_dir <<< blocks, threads, share_bytes >>>(gpuDir, texObj, gpuGrad, gpuTable, phi, rmax, x , y);
//call the kernel to update the gradient direction
cuda_update_grad <<< blocks, threads >>>(gpuGrad, gpuDir, x , y);
//free allocated memory
cudaFree(gpuDir);
}
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
|
