3f0de7dd
Laila Saadatifard
upload the vote a...
|
1
2
|
#ifndef STIM_CUDA_UPDATE_DIR_THRESHOLD_GLOBALD_H
#define STIM_CUDA_UPDATE_DIR_THRESHOLD_GLOBAL_H
|
03428452
Laila Saadatifard
update the ivote ...
|
3
4
5
6
7
|
# include <iostream>
# include <cuda.h>
#include <stim/cuda/cudatools.h>
#include <stim/cuda/sharedmem.cuh>
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
8
|
#include "cpyToshare.cuh"
|
03428452
Laila Saadatifard
update the ivote ...
|
9
10
11
12
13
14
|
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>
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
15
|
__global__ void cuda_update_dir(T* gpuDir, T* gpuVote, T* gpuTh, T* gpuTable, T phi, int rmax, int th_size, int x, int y){
|
03428452
Laila Saadatifard
update the ivote ...
|
16
|
|
03428452
Laila Saadatifard
update the ivote ...
|
17
|
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
18
19
|
// calculate the coordinate for this current thread.
|
03428452
Laila Saadatifard
update the ivote ...
|
20
|
int xi = blockIdx.x * blockDim.x + threadIdx.x;
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
21
22
|
// calculate the voting direction based on the grtadient direction
float theta = gpuTh[3*xi];
|
03428452
Laila Saadatifard
update the ivote ...
|
23
|
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
24
25
26
27
|
//calculate the position and x, y coordinations of this voter in the original image
unsigned int i_v = gpuTh[3*xi+2];
unsigned int y_v = i_v/x;
unsigned int x_v = i_v - (y_v*x);
|
03428452
Laila Saadatifard
update the ivote ...
|
28
|
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
29
30
31
32
33
34
35
36
37
38
39
40
|
//initialize the vote direction to zero
gpuDir[xi] = 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
int x_table = 2*rmax +1;
|
03428452
Laila Saadatifard
update the ivote ...
|
41
42
|
int rmax_sq = rmax * rmax;
int tx_rmax = threadIdx.x + rmax;
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
|
if(xi < th_size){
for(int yr = -rmax; yr <= rmax; yr++){
for(int xr = -rmax; xr <= rmax; xr++){
unsigned int ind_t = (rmax - yr) * x_table + rmax - xr;
// find the angle between the voter and the current pixel in x and y directions
float atan_angle = gpuTable[ind_t];
// 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)){
// find the vote value for the current counter
float vote_c = gpuVote[(y_v+yr)*x + (x_v+xr)];
// 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;
|
03428452
Laila Saadatifard
update the ivote ...
|
64
65
66
67
|
}
}
}
}
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
68
|
|
03428452
Laila Saadatifard
update the ivote ...
|
69
|
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
70
71
72
73
|
unsigned int ind_m = (rmax - id_y) * x_table + (rmax - id_x);
float new_angle = gpuTable[ind_m];
gpuDir[xi] = new_angle;
}
|
03428452
Laila Saadatifard
update the ivote ...
|
74
|
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
75
|
}
|
03428452
Laila Saadatifard
update the ivote ...
|
76
77
78
|
// this kernel updates the gradient direction by the calculated voting direction.
template<typename T>
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
79
|
__global__ void cuda_update_grad(T* gpuTh, T* gpuDir, int th_size, int x, int y){
|
03428452
Laila Saadatifard
update the ivote ...
|
80
|
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
81
|
// calculate the coordinate for this current thread.
|
03428452
Laila Saadatifard
update the ivote ...
|
82
|
int xi = blockIdx.x * blockDim.x + threadIdx.x;
|
03428452
Laila Saadatifard
update the ivote ...
|
83
|
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
84
|
|
03428452
Laila Saadatifard
update the ivote ...
|
85
|
//update the gradient image with the vote direction
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
86
|
gpuTh[3*xi] = gpuDir[xi];
|
03428452
Laila Saadatifard
update the ivote ...
|
87
88
89
|
}
template<typename T>
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
90
|
void gpu_update_dir(T* gpuVote, T* gpuTh, T* gpuTable, T phi, unsigned int rmax, unsigned int th_size, unsigned int x, unsigned int y){
|
03428452
Laila Saadatifard
update the ivote ...
|
91
92
|
//calculate the number of bytes in the array
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
93
|
unsigned int bytes_th = th_size* sizeof(T);
|
03428452
Laila Saadatifard
update the ivote ...
|
94
95
|
unsigned int max_threads = stim::maxThreadsPerBlock();
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
96
97
|
dim3 threads(max_threads);
dim3 blocks(th_size/threads.x+1);
|
03428452
Laila Saadatifard
update the ivote ...
|
98
99
100
|
// allocate space on the GPU for the updated vote direction
T* gpuDir;
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
101
|
cudaMalloc(&gpuDir, bytes_th);
|
03428452
Laila Saadatifard
update the ivote ...
|
102
103
|
//call the kernel to calculate the new voting direction
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
104
|
cuda_update_dir <<< blocks, threads>>>(gpuDir, gpuVote, gpuTh, gpuTable, phi, rmax, th_size, x , y);
|
03428452
Laila Saadatifard
update the ivote ...
|
105
106
|
//call the kernel to update the gradient direction
|
3f0de7dd
Laila Saadatifard
upload the vote a...
|
107
|
cuda_update_grad <<< blocks, threads >>>(gpuTh, gpuDir, th_size, x , y);
|
03428452
Laila Saadatifard
update the ivote ...
|
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
|
//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
|