vote.cuh
5.84 KB
1
2
3
4
5
6
7
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
#ifndef STIM_CUDA_VOTE_H
#define STIM_CUDA_VOTE_H
# include <iostream>
# include <cuda.h>
#include <stim/cuda/cudatools.h>
#include <stim/cuda/sharedmem.cuh>
namespace stim{
namespace cuda{
// this kernel calculates the vote value by adding up the gradient magnitudes of every voter that this pixel is located in their voting area
template<typename T>
__global__ void cuda_vote(T* gpuVote, cudaTextureObject_t in, T* gpuTable, T phi, int rmax, int x, int y){
//generate a pointer to shared memory (size will be specified as a kernel parameter)
extern __shared__ float2 s_grad[];
//calculate the start point for this block
int bxi = blockIdx.x * blockDim.x;
// calculate the 2D coordinates for this current thread.
int xi = bxi + threadIdx.x;
int yi = blockIdx.y * blockDim.y + threadIdx.y;
// convert 2D coordinates to 1D
int i = yi * x + xi;
// define a local variable to sum the votes from the voters
float sum = 0;
//calculate the width of the shared memory block
int swidth = 2 * rmax + blockDim.x;
// compute the size of window which will be checked for finding the proper voters for this pixel
int x_table = 2*rmax +1;
int rmax_sq = rmax * rmax;
int tx_rmax = threadIdx.x + rmax;
int bxs = bxi - rmax;
//for every line (along y)
for(int yr = -rmax; yr <= rmax; yr++){
//copy the portion of the image necessary for this block to shared memory
__syncthreads();
stim::cuda::sharedMemcpy_tex2D<float2>(s_grad, in, bxs, yi + yr , swidth, 1, threadIdx, blockDim);
__syncthreads();
if(xi < x && yi < y){
for(int xr = -rmax; xr <= rmax; xr++){
//find the location of this voter in the atan2 table
int id_t = (yr + rmax) * x_table + xr + rmax;
// calculate the angle between the pixel and the current voter in x and y directions
float atan_angle = gpuTable[id_t];
// calculate the voting direction based on the grtadient direction
int idx_share = xr + tx_rmax ;
float2 g = s_grad[idx_share];
float theta = g.x;
// check if the current voter is located in the voting area of this pixel.
if (((xr * xr + yr *yr)< rmax_sq) && (abs(atan_angle - theta) <phi)){
sum += g.y;
}
}
}
}
if(xi < x && yi < y)
gpuVote[i] = sum;
}
template<typename T>
void gpu_vote(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, 32, 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, gpuGrad, bytes*2,
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)*2*4;
//call the kernel to do the voting
cuda_vote <<< blocks, threads,share_bytes >>>(gpuVote, texObj, gpuTable, phi, rmax, x , y);
cudaDestroyTextureObject(texObj);
cudaFreeArray(cuArray);
}
template<typename T>
void cpu_vote(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);
//allocate space on the GPU for the input Gradient image
T* gpuGrad;
HANDLE_ERROR(cudaMalloc(&gpuGrad, bytes*2));
//copy the Gradient Magnitude 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 vote calculation function
gpu_vote<T>(gpuVote, gpuGrad, gpuTable, phi, rmax, x , y);
//copy the Vote Data back to the CPU
cudaMemcpy(cpuVote, gpuVote, bytes, cudaMemcpyDeviceToHost) ;
//free allocated memory
cudaFree(gpuTable);
cudaFree(gpuVote);
cudaFree(gpuGrad);
}
}
}
#endif