6ef1dab9
Laila Saadatifard
fix one bug in th...
|
1
2
3
|
#include "cuda_fp16.h"
#include "float_to_half.cuh"
#include "half_to_float.cuh"
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
4
5
|
#include "gaussian_blur3.cuh"
#include "gradient3.cuh"
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
6
7
8
9
|
#include "mag3.cuh"
#include "vote3.cuh"
#include "update_dir3.cuh"
#include "local_max3.cuh"
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
10
11
|
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
12
13
14
|
void ivote3(float* center, float* img, float sigma[], float phi, float d_phi, unsigned int r[],
int iter, float t, unsigned int conn[], unsigned int x, unsigned int y, unsigned int z){
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
15
|
// compute the number of bytes in the input data
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
16
|
unsigned int bytes = x * y * z * sizeof(float);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
17
18
|
//assign memory on gpu for the input data.
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
19
|
float* gpuI0;
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
20
21
22
|
cudaMalloc(&gpuI0, bytes);
//copy the image data to the GPU.
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
23
24
|
cudaMemcpy(gpuI0, img, bytes, cudaMemcpyHostToDevice);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
25
|
//call the blurring function from the gpu.
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
26
27
|
gpu_gaussian_blur3<float>(gpuI0, sigma, x, y, z);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
28
29
|
cudaDeviceSynchronize();
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
30
|
//assign memory on the gpu for the gradient along the X, y, z.
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
31
32
33
|
float* gpu_grad;
cudaMalloc(&gpu_grad, bytes*3);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
34
35
36
37
|
//call the gradient function from the gpu.
gpu_gradient3<float>(gpu_grad, gpuI0, x, y, z);
cudaFree(gpuI0);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
38
39
|
float* gpu_vote;
cudaMalloc(&gpu_vote, bytes);
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
40
|
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
41
42
43
44
45
|
float cos_phi = cos(phi);
//call the vote function.
for (int i = 0; i < iter; i++){
|
02fb26b3
Laila Saadatifard
change the vote a...
|
46
|
gpu_vote3<float>(gpu_vote, gpu_grad, cos_phi, r, x, y, z);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
47
48
49
50
51
52
53
54
55
56
57
58
|
cudaDeviceSynchronize();
if (i==0)
cudaMemcpy(img, gpu_vote, bytes, cudaMemcpyDeviceToHost);
if (phi >= d_phi){
gpu_update_dir3<float>(gpu_grad, gpu_vote, cos_phi, r, x, y, z);
cudaDeviceSynchronize();
phi = phi - d_phi;
cos_phi = cos(phi);
}
}
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
59
|
|
02fb26b3
Laila Saadatifard
change the vote a...
|
60
|
cudaFree(gpu_grad);
|
6ef1dab9
Laila Saadatifard
fix one bug in th...
|
61
|
cudaMemcpy(center, gpu_vote, bytes, cudaMemcpyDeviceToHost);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
62
63
64
65
66
67
68
69
70
|
//allocate space on the gpu for the final detected cells.
float* gpu_output;
cudaMalloc(&gpu_output, bytes);
//call the local max function
gpu_local_max3<float>(gpu_output, gpu_vote, t, conn, x, y, z);
//copy the final result to the cpu.
|
6ef1dab9
Laila Saadatifard
fix one bug in th...
|
71
|
//cudaMemcpy(center, gpu_output, bytes, cudaMemcpyDeviceToHost);
|
f12505fb
Laila Saadatifard
upload the ivote ...
|
72
73
74
75
76
|
cudaFree(gpu_vote);
cudaFree(gpu_output);
|
c9aba18a
Laila Saadatifard
ivote3 on the GPU...
|
77
|
}
|