5c079506
Laila Saadatifard
upload the ivote ...
|
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
|
// 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 update_dir3(T* gpu_dir, T* gpu_grad, T* gpu_vote, T cos_phi, int rx, int ry, int rz, int x, int y, int z){
//calculate x,y,z coordinates for this thread
int xi = blockIdx.x * blockDim.x + threadIdx.x;
//find the grid size along y
int grid_y = y / blockDim.y;
int blockidx_y = blockIdx.y % grid_y;
int yi = blockidx_y * blockDim.y + threadIdx.y;
int zi = blockIdx.y / grid_y;
int i = zi * x * y + yi * x + xi;
if(xi >= x|| yi >= y || zi>= z) return;
//find the gradient values along the x, y ,z axis, and the gradient magnitude for the voter
float g_v_x = gpu_grad[i * 3 + 0];
float g_v_y = gpu_grad[i * 3 + 1];
float g_v_z = gpu_grad[i * 3 + 2];
|
5c079506
Laila Saadatifard
upload the ivote ...
|
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
|
}
// this kernel updates the gradient direction by the calculated voting direction.
template<typename T>
__global__ void update_grad3(T* gpu_grad, T* gpu_dir, int x, int y, int z){
//calculate x,y,z coordinates for this thread
int xi = blockIdx.x * blockDim.x + threadIdx.x;
//find the grid size along y
int grid_y = y / blockDim.y;
int blockidx_y = blockIdx.y % grid_y;
int yi = blockidx_y * blockDim.y + threadIdx.y;
int zi = blockIdx.y / grid_y;
int i = zi * x * y + yi * x + xi;
if(xi >= x || yi >= y || zi >= z) return;
//update the gradient image with the new direction direction
gpu_grad[i * 3 + 0] = gpu_dir [i * 3 + 0];
gpu_grad[i * 3 + 1] = gpu_dir [i * 3 + 1];
gpu_grad[i * 3 + 2] = gpu_dir [i * 3 + 2];
}
template<typename T>
void gpu_update_dir3(T* gpu_grad, T* gpu_vote, T cos_phi, unsigned int r[], unsigned int x, unsigned int y, unsigned int z){
unsigned int max_threads = stim::maxThreadsPerBlock();
dim3 threads(sqrt (max_threads),sqrt (max_threads));
dim3 blocks(x / threads.x + 1, (y / threads.y + 1) * z);
// allocate space on the GPU for the updated vote direction
T* gpu_dir;
cudaMalloc(&gpu_dir, x * y * z * sizeof(T) * 3);
//call the kernel to calculate the new voting direction
update_dir3 <<< blocks, threads >>>(gpu_dir, gpu_grad, gpu_vote, cos_phi, r[0], r[1], r[2], x , y, z);
//call the kernel to update the gradient direction
update_grad3 <<< blocks, threads >>>(gpu_grad, gpu_dir, x , y, z);
//free allocated memory
cudaFree(gpu_dir);
}
template<typename T>
void cpu_update_dir3(T* cpu_grad, T* cpu_vote, T cos_phi, unsigned int r[], unsigned int x, unsigned int y, unsigned int z){
//calculate the number of bytes in the array
unsigned int bytes = x * y * z * sizeof(T);
//allocate space on the GPU for the Vote data
T* gpu_vote;
cudaMalloc(&gpu_vote, bytes);
//copy the input vote data to the GPU
cudaMemcpy(gpu_vote, cpu_vote, bytes, cudaMemcpyHostToDevice);
//allocate space on the GPU for the Gradient data
T* gpu_grad;
cudaMalloc(&gpu_grad, bytes*3);
//copy the Gradient data to the GPU
cudaMemcpy(gpu_grad, cpu_grad, bytes*3, cudaMemcpyHostToDevice);
//call the GPU version of the update direction function
gpu_update_dir3<T>(gpu_grad, gpu_vote, cos_phi, r, x , y, z);
//copy the new gradient image back to the CPU
cudaMemcpy(cpu_grad, gpu_grad, bytes*3, cudaMemcpyDeviceToHost) ;
//free allocated memory
cudaFree(gpu_vote);
cudaFree(gpu_grad);
}
#endif
|