13fe3c84
Laila Saadatifard
update the stimli...
|
1
2
3
|
#ifndef STIM_CUDA_VOTE_H
#define STIM_CUDA_VOTE_H
|
13fe3c84
Laila Saadatifard
update the stimli...
|
4
5
|
# include <iostream>
# include <cuda.h>
|
96f9b10f
Laila Saadatifard
change the header...
|
6
|
#include <stim/cuda/cudatools.h>
|
13fe3c84
Laila Saadatifard
update the stimli...
|
7
8
|
#include <stim/cuda/sharedmem.cuh>
|
13fe3c84
Laila Saadatifard
update the stimli...
|
9
10
11
12
13
|
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>
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
14
|
__global__ void cuda_vote(T* gpuVote, cudaTextureObject_t in, T* gpuTable, T phi, int rmax, int x, int y){
|
13fe3c84
Laila Saadatifard
update the stimli...
|
15
16
17
18
19
20
21
|
//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;
|
13fe3c84
Laila Saadatifard
update the stimli...
|
22
23
|
// calculate the 2D coordinates for this current thread.
int xi = bxi + threadIdx.x;
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
24
|
int yi = blockIdx.y * blockDim.y + threadIdx.y;
|
13fe3c84
Laila Saadatifard
update the stimli...
|
25
26
|
// convert 2D coordinates to 1D
int i = yi * x + xi;
|
bf731970
Laila Saadatifard
fix some bugs in ...
|
27
|
|
13fe3c84
Laila Saadatifard
update the stimli...
|
28
29
|
// define a local variable to sum the votes from the voters
float sum = 0;
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
30
|
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
31
32
33
|
//calculate the width of the shared memory block
int swidth = 2 * rmax + blockDim.x;
|
13fe3c84
Laila Saadatifard
update the stimli...
|
34
|
// compute the size of window which will be checked for finding the proper voters for this pixel
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
35
36
|
int x_table = 2*rmax +1;
int rmax_sq = rmax * rmax;
|
13fe3c84
Laila Saadatifard
update the stimli...
|
37
38
39
|
int tx_rmax = threadIdx.x + rmax;
int bxs = bxi - rmax;
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
40
41
|
//for every line (along y)
for(int yr = -rmax; yr <= rmax; yr++){
|
13fe3c84
Laila Saadatifard
update the stimli...
|
42
43
44
45
46
47
|
//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();
|
13fe3c84
Laila Saadatifard
update the stimli...
|
48
|
if(xi < x && yi < y){
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
49
50
|
for(int xr = -rmax; xr <= rmax; xr++){
|
13fe3c84
Laila Saadatifard
update the stimli...
|
51
52
|
//find the location of this voter in the atan2 table
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
53
|
int id_t = (yr + rmax) * x_table + xr + rmax;
|
13fe3c84
Laila Saadatifard
update the stimli...
|
54
55
56
|
// calculate the angle between the pixel and the current voter in x and y directions
float atan_angle = gpuTable[id_t];
|
bf731970
Laila Saadatifard
fix some bugs in ...
|
57
|
|
13fe3c84
Laila Saadatifard
update the stimli...
|
58
59
60
61
62
63
64
65
66
67
|
// 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;
}
|
13fe3c84
Laila Saadatifard
update the stimli...
|
68
|
}
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
69
|
|
13fe3c84
Laila Saadatifard
update the stimli...
|
70
71
|
}
}
|
84ca9bba
Laila Saadatifard
fix some bugs in ...
|
72
73
74
|
if(xi < x && yi < y)
gpuVote[i] = sum;
|
13fe3c84
Laila Saadatifard
update the stimli...
|
75
76
77
78
79
80
81
82
|
}
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;
|
bf731970
Laila Saadatifard
fix some bugs in ...
|
83
|
|
13fe3c84
Laila Saadatifard
update the stimli...
|
84
|
unsigned int max_threads = stim::maxThreadsPerBlock();
|
13fe3c84
Laila Saadatifard
update the stimli...
|
85
86
|
dim3 threads(max_threads, 1);
dim3 blocks(x/threads.x + (x %threads.x == 0 ? 0:1) , y);
|
bf731970
Laila Saadatifard
fix some bugs in ...
|
87
|
|
13fe3c84
Laila Saadatifard
update the stimli...
|
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
|
// 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;
|
13fe3c84
Laila Saadatifard
update the stimli...
|
123
124
125
126
|
//call the kernel to do the voting
cuda_vote <<< blocks, threads,share_bytes >>>(gpuVote, texObj, gpuTable, phi, rmax, x , y);
|
59781ee3
Pavel Govyadinov
fixed a stask bug...
|
127
128
129
|
cudaDestroyTextureObject(texObj);
cudaFreeArray(cuArray);
|
13fe3c84
Laila Saadatifard
update the stimli...
|
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
|
}
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));
|
bf731970
Laila Saadatifard
fix some bugs in ...
|
159
|
|
13fe3c84
Laila Saadatifard
update the stimli...
|
160
161
|
//call the GPU version of the vote calculation function
gpu_vote<T>(gpuVote, gpuGrad, gpuTable, phi, rmax, x , y);
|
bf731970
Laila Saadatifard
fix some bugs in ...
|
162
|
|
13fe3c84
Laila Saadatifard
update the stimli...
|
163
164
165
166
167
168
169
170
171
172
173
174
|
//copy the Vote Data back to the CPU
cudaMemcpy(cpuVote, gpuVote, bytes, cudaMemcpyDeviceToHost) ;
//free allocated memory
cudaFree(gpuTable);
cudaFree(gpuVote);
cudaFree(gpuGrad);
}
}
}
|
59781ee3
Pavel Govyadinov
fixed a stask bug...
|
175
|
#endif
|