Commit 94d437dd6ed950af9ba7b5f1d3596800697adccf
1 parent
02fb26b3
ivote3 code compiling on the gpu and use three channels for voting direction and gradient magnitude
Showing
8 changed files
with
81 additions
and
506 deletions
Show diff stats
Matlab_3D/main.m deleted
1 | - | |
2 | -clc; | |
3 | -clear; | |
4 | -disp('***************** NEW RUN *********************'); | |
5 | -total = tic; | |
6 | - | |
7 | - | |
8 | -% ******* Initialize voting parameters ************************************** | |
9 | -rmax = 9; %maximum radius of the cell | |
10 | -ang_deg = 15.1; %half the angular range of the voting area | |
11 | -ang = ang_deg * pi / 180; | |
12 | -iter = 6; %number of voting iterations | |
13 | -t0 = 1.0; %threshold color | |
14 | -sigma = [2, 2, 1]; | |
15 | -% t = 0.1; | |
16 | -iter = iter-1; | |
17 | -d_ang= ang / (iter); | |
18 | - | |
19 | -% ******** Testing parameters ****************************************** | |
20 | -p = [100, 50, 100]; | |
21 | -ps = [100, 100, 50]; | |
22 | -% ps = [100, 50, 40]; | |
23 | -% I = syn_Img(rmax , ps); | |
24 | -volfile = 'nissl-rat.vol'; | |
25 | -fid = fopen(volfile); % open the file that include the image | |
26 | -S = fread(fid, 3, 'int32'); | |
27 | -X = S(1); | |
28 | -Y = S(2); | |
29 | -Z = S(3); | |
30 | - | |
31 | -% load the VOL data into a 2D matrix | |
32 | -I = fread(fid,[X Y*Z], 'uint8'); | |
33 | -fclose(fid); | |
34 | - | |
35 | -%change this to a 3D matrix | |
36 | -I = (reshape(I, [X, Y, Z])); | |
37 | - | |
38 | -% invert the intensity | |
39 | -I = 255 - I; | |
40 | - | |
41 | -%perform a gaussian blur | |
42 | -Iblur = gauss_blur3d(I, sigma); | |
43 | - | |
44 | -%crop out a small subregion of I and Iblur | |
45 | -Iblur = Iblur(p(1):p(1)+ps(1)-1, p(2):p(2)+ps(2)-1, p(3):p(3)+ps(3)-1); | |
46 | -I = I(p(1):p(1)+ps(1)-1, p(2):p(2)+ps(2)-1, p(3):p(3)+ps(3)-1); | |
47 | - | |
48 | -% compute the gradient | |
49 | -[Igrad_y, Igrad_x, Igrad_z] = gradient(Iblur); | |
50 | - | |
51 | -%calculate the gradient magnitude | |
52 | -Imag = sqrt(Igrad_x .^ 2 + Igrad_y .^ 2 + Igrad_z .^2); | |
53 | - | |
54 | -%set a threshold for the gradient magnitude | |
55 | -It = Imag > t0; | |
56 | - | |
57 | -%Set the boundaries of the threshold image to zero | |
58 | -It(1:rmax, :, :) = 0; | |
59 | -It(ps(1) - rmax:ps(1), :,:) = 0; | |
60 | -It(:, 1:rmax, :) = 0; | |
61 | -It(:, ps(2) - rmax:ps(2),:) = 0; | |
62 | -It(:, :, 1:rmax) = 0; | |
63 | -It(:,:, ps(3) - rmax:ps(3)) = 0; | |
64 | - | |
65 | -%get the indices of all of the nonzero values in the threshold image | |
66 | -% (voter positions) | |
67 | -[Itx,Ity,Itz] = ind2sub(size(It),find(It)); | |
68 | -Vi = find(It); | |
69 | -nV = nnz(It); | |
70 | - | |
71 | -% create a meshgrid describing coordinates relative to the voter position | |
72 | -range = -rmax:rmax; %create an array of values between -rmax and rmax | |
73 | -[mx, my, mz] = meshgrid(range, range, range); %create a template describing local pixel position in a small cube | |
74 | -m_mag = sqrt(mx.^2 + my.^2 + mz.^2); %create a template describing the distance from the center of a small cube | |
75 | - | |
76 | -% create a mask for the voting area | |
77 | -M_dist = m_mag <= rmax; %mask for the voting area distance (all values < rmax from the center) | |
78 | - | |
79 | -%calculate the direction vector between a pixel and voter | |
80 | -LV_x = mx./m_mag; | |
81 | -LV_y = my./m_mag; | |
82 | -LV_z = mz./m_mag; | |
83 | - | |
84 | -%number of pixels in the voting area of each voter (initialize to zero) | |
85 | -validPixels = (zeros(nV,1)); | |
86 | - | |
87 | -%indices of pixels in the voting area of each voter | |
88 | -% indices reference the 3D image | |
89 | -g_v_prime = (zeros(nV, (rmax^3))); | |
90 | - | |
91 | - | |
92 | -%% vote | |
93 | -tic; | |
94 | -%for each iteration (in iterative voting) | |
95 | -for itr = 1 : iter+1 | |
96 | - | |
97 | - %initialize the vote image to zero | |
98 | - Ivote = (zeros(size(I))); | |
99 | - | |
100 | - %for each voter (nonzero pixels in the threshold image It) | |
101 | - for v = 1: nV | |
102 | - | |
103 | - %get the cartesian coordinates of the voter v in the main image I | |
104 | - vx = Itx(v); | |
105 | - vy = Ity(v); | |
106 | - vz = Itz(v); | |
107 | - vi = Vi(v); | |
108 | - | |
109 | - %retreive the gradient magnitude at the voter position | |
110 | - %vmag = Imag(vx,vy,vz); | |
111 | - vmag = Imag(vi); | |
112 | - | |
113 | - %retrieve the gradient | |
114 | - gx = Igrad_x(vi); | |
115 | - gy = Igrad_y(vi); | |
116 | - gz = Igrad_z(vi); | |
117 | - | |
118 | - %calculate the gradient magnitude | |
119 | - dmag = sqrt (gx^2 + gy^2 + gz^2); | |
120 | - | |
121 | - %calculate the normalized gradient direction | |
122 | - dx = gx / dmag; | |
123 | - dy = gy / dmag; | |
124 | - dz = gz / dmag; | |
125 | - | |
126 | - %calculate the angle between the voter direction and the pixel direction | |
127 | - cos_diff = LV_x .* dx + LV_y .* dy + LV_z .* dz; | |
128 | - ang_diff = acos(cos_diff); | |
129 | - | |
130 | - %create an angular mask for the voting area | |
131 | - M_angle = cos_diff > cos(ang); | |
132 | - | |
133 | - %combine the two masks to mask out the voting angle | |
134 | - M = M_angle .* M_dist; | |
135 | - | |
136 | - % get the coordinates of each pixel in the final voter mask M | |
137 | - pi = find(M); | |
138 | - | |
139 | - %calculate the number of pixels in the voting region | |
140 | - npts = nnz(M); | |
141 | - validPixels(v) = npts; | |
142 | - | |
143 | - %convert every index in the voting area from a local 3D index to a global 3D index (into the original image I) | |
144 | - global_px = vx + mx(pi); | |
145 | - global_py = vy + my(pi); | |
146 | - global_pz = vz + mz(pi); | |
147 | - | |
148 | - %convert the global 3D index of each point into a global 1D index | |
149 | - global_pi = sub2ind(ps, global_px, global_py, global_pz); | |
150 | - | |
151 | - g_v_prime (v, 1:npts) = global_pi; | |
152 | - | |
153 | - | |
154 | - Ivote( global_pi ) = Ivote( global_pi ) + vmag; | |
155 | - | |
156 | - end | |
157 | - | |
158 | - if itr ==1 | |
159 | - Ivote1 = Ivote; | |
160 | - | |
161 | - elseif itr ==2 | |
162 | - Ivote2 = Ivote; | |
163 | - | |
164 | - elseif itr ==3 | |
165 | - Ivote3 = Ivote; | |
166 | - | |
167 | - elseif itr ==4 | |
168 | - Ivote4 = Ivote; | |
169 | - | |
170 | - elseif itr == 5 | |
171 | - Ivote5 = Ivote; | |
172 | - end | |
173 | - t_v1 = toc; | |
174 | - disp(['voting done. time =',num2str(t_v1)]); | |
175 | - | |
176 | - % update the voting direction | |
177 | - if ang>0 | |
178 | - tic; | |
179 | - for v = 1: nV | |
180 | - % coordinates of the current voter | |
181 | - vx = Itx(v); | |
182 | - vy = Ity(v); | |
183 | - vz = Itz(v); | |
184 | - | |
185 | - %get the local value of the voting image | |
186 | - local_Ivote = Ivote(g_v_prime(v,1:validPixels(v))); | |
187 | - | |
188 | - %find the index of the maximum value | |
189 | - [~, local_max_idx] = max(local_Ivote); | |
190 | - | |
191 | - %convert this into a global subscript | |
192 | - [g_px, g_py, g_pz] = ind2sub(size(Ivote), g_v_prime(v,local_max_idx)); | |
193 | - | |
194 | - %compute the vector from the voter position to this position | |
195 | - Igrad_x(vx, vy, vz) = g_px - vx; | |
196 | - Igrad_y(vx, vy, vz) = g_py - vy; | |
197 | - Igrad_z(vx, vy, vz) = g_pz - vz; | |
198 | - | |
199 | - end | |
200 | - | |
201 | - | |
202 | - tdir1 = toc; | |
203 | - display (['updating dir done. time = ', num2str(tdir1)]); | |
204 | - ang = ang - d_ang; | |
205 | - end | |
206 | - end | |
207 | - | |
208 | - | |
209 | -%% | |
210 | -% t = 50; | |
211 | -% center = Ivote; | |
212 | -% % center(center<t) = 0; | |
213 | -% % center = imregionalmax(center); | |
214 | -% cn = nnz(center); | |
215 | -% [cx, cy, cz] = ind2sub(size(center), find(center)); | |
216 | -% c2d = zeros(size(center,1), size(center,2)); | |
217 | -% for cc =1:cn | |
218 | -% c2d(cx(cc), cy(cc)) = 1; | |
219 | -% end | |
220 | -% I2d = max(Iblur, [], 3); | |
221 | -% % figure(1),imagesc(I2d); colormap(gray); | |
222 | -% figure(2),imagesc(c2d); colormap(gray); | |
223 | -% | |
224 | -% % out(:,:,1) = mat2gray(I2d); | |
225 | -% % out(:,:,2) = mat2gray(c2d); | |
226 | -% % out(:,:,3) = mat2gray(I2d); | |
227 | -% % figure (5), imagesc(out); | |
228 | -% imwrite(mat2gray(c2d), 'vote.bmp'); | |
229 | -figure(6); imagesc(squeeze(Ivote1(:,ceil(size(Ivote1,2)/2),:))); colormap(gray); |
Matlab_3D/main_rmax.m deleted
1 | - | |
2 | - | |
3 | -clc; | |
4 | -clear; | |
5 | -disp('***************** NEW RUN *********************'); | |
6 | -total = tic; | |
7 | - | |
8 | - | |
9 | -% ******* Initialize voting parameters ************************************** | |
10 | -rmax = [9 9 5]; %maximum radius of the cell | |
11 | -ang_deg = 20.1; %half the angular range of the voting area | |
12 | -ang = ang_deg * pi / 180; | |
13 | -iter = 5; %number of voting iterations | |
14 | -t0 = 1.0; %threshold color | |
15 | -sigma = [3, 3, 1.5]; | |
16 | -% t = 0.1; | |
17 | - | |
18 | -d_ang= ang / (iter); | |
19 | - | |
20 | -% ******** Testing parameters ****************************************** | |
21 | -p = [50, 50, 150]; | |
22 | -ps = [400, 400, 200]; | |
23 | -% ps = [100, 50, 40]; | |
24 | -% I = syn_Img(rmax , ps); | |
25 | -volfile = 'nissl-rat.vol'; | |
26 | -fid = fopen(volfile); % open the file that include the image | |
27 | -S = fread(fid, 3, 'int32'); | |
28 | -X = S(1); | |
29 | -Y = S(2); | |
30 | -Z = S(3); | |
31 | -% load the VOL data into a 2D matrix | |
32 | -I = fread(fid,[X Y*Z], 'uint8'); | |
33 | -fclose(fid); | |
34 | -%change this to a 3D matrix | |
35 | -I = (reshape(I, [X, Y, Z])); | |
36 | -% invert the intensity | |
37 | -I = (255 - I); | |
38 | - | |
39 | -%perform a gaussian blur | |
40 | -Iblur = gauss_blur3d(I, sigma); | |
41 | - | |
42 | -%crop out a small subregion of I and Iblur | |
43 | -Iblur = Iblur(p(1):p(1)+ps(1)-1, p(2):p(2)+ps(2)-1, p(3):p(3)+ps(3)-1); | |
44 | -I = I(p(1):p(1)+ps(1)-1, p(2):p(2)+ps(2)-1, p(3):p(3)+ps(3)-1); | |
45 | -% | |
46 | -% compute the gradient | |
47 | -[Igrad_y, Igrad_x, Igrad_z] = gradient(Iblur); | |
48 | - | |
49 | -%calculate the gradient magnitude | |
50 | -Imag = sqrt(Igrad_x .^ 2 + Igrad_y .^ 2 + Igrad_z .^2); | |
51 | -Isize = size(I); | |
52 | -I = single(I); | |
53 | -Iblur = single(Iblur); | |
54 | - | |
55 | -% | |
56 | -%set a threshold for the gradient magnitude | |
57 | -It = Imag > t0; | |
58 | - | |
59 | -%Set the boundaries of the threshold image to zero | |
60 | -It(1:rmax(1), :, :) = 0; | |
61 | -It(ps(1) - rmax(1):ps(1), :,:) = 0; | |
62 | -It(:, 1:rmax(2), :) = 0; | |
63 | -It(:, ps(2) - rmax(2):ps(2),:) = 0; | |
64 | -It(:, :, 1:rmax(3)) = 0; | |
65 | -It(:,:, ps(3) - rmax(3):ps(3)) = 0; | |
66 | -% | |
67 | -%get the indices of all of the nonzero values in the threshold image | |
68 | -% (voter positions) | |
69 | -[Itx,Ity,Itz] = ind2sub(size(It),find(It)); | |
70 | -Vi =(find(It)); | |
71 | -nV = nnz(It); | |
72 | -% | |
73 | -% create a meshgrid describing coordinates relative to the voter position | |
74 | -rangex = -rmax(1):rmax(1); %create an array of values between -rmax and rmax | |
75 | -rangey = -rmax(2):rmax(2); | |
76 | -rangez = -rmax(3):rmax(3); | |
77 | -[mx, my, mz] = meshgrid(rangex, rangey, rangez); %create a template describing local pixel position in a small cube | |
78 | -m_mag = (sqrt(mx.^2 + my.^2 + mz.^2)); %create a template describing the distance from the center of a small cube | |
79 | - | |
80 | -% create a mask for the voting area | |
81 | -M_dist = (mx.^2/rmax(1)^2 + my.^2/rmax(2)^2 + mz.^2/rmax(3)^2) <= 1; %mask for the voting area distance (all values < rmax from the center) | |
82 | - | |
83 | -% calculate the direction vector between a pixel and voter | |
84 | -LV_x = mx./m_mag; | |
85 | -LV_y = my./m_mag; | |
86 | -LV_z = mz./m_mag; | |
87 | - | |
88 | -%number of pixels in the voting area of each voter (initialize to zero) | |
89 | -validPixels = (zeros(nV,1)); | |
90 | -%% | |
91 | -%indices of pixels in the voting area of each voter | |
92 | -% indices reference the 3D image | |
93 | -g_v_prime = zeros(nV, ceil(rmax(1)*rmax(2)*rmax(3)*ang)); | |
94 | - | |
95 | - | |
96 | -%% vote | |
97 | -tic; | |
98 | -%for each iteration (in iterative voting) | |
99 | -for itr = 1 : iter+1 | |
100 | - | |
101 | - %initialize the vote image to zero | |
102 | - Ivote = zeros(Isize); | |
103 | - | |
104 | - %for each voter (nonzero pixels in the threshold image It) | |
105 | - for v = 1: nV | |
106 | - | |
107 | - %get the cartesian coordinates of the voter v in the main image I | |
108 | - vx = Itx(v); | |
109 | - vy = Ity(v); | |
110 | - vz = Itz(v); | |
111 | - vi = Vi(v); | |
112 | - | |
113 | - %retreive the gradient magnitude at the voter position | |
114 | - vmag = Imag(vi); | |
115 | - | |
116 | - %retrieve the gradient | |
117 | - gx = Igrad_x(vi); | |
118 | - gy = Igrad_y(vi); | |
119 | - gz = Igrad_z(vi); | |
120 | - | |
121 | - %calculate the gradient magnitude | |
122 | - dmag = sqrt (gx^2 + gy^2 + gz^2); | |
123 | - | |
124 | - %calculate the normalized gradient direction | |
125 | - dx = gx / dmag; | |
126 | - dy = gy / dmag; | |
127 | - dz = gz / dmag; | |
128 | - | |
129 | - %calculate the angle between the voter direction and the pixel direction | |
130 | - cos_diff = LV_x .* dx + LV_y .* dy + LV_z .* dz; | |
131 | - ang_diff = acos(cos_diff); | |
132 | - | |
133 | - %create an angular mask for the voting area | |
134 | - M_angle = cos_diff >= cos(ang); | |
135 | - | |
136 | - %combine the two masks to mask out the voting angle | |
137 | - M = M_angle .* M_dist; | |
138 | - | |
139 | - % get the coordinates of each pixel in the final voter mask M | |
140 | - pi = find(M); | |
141 | - | |
142 | - %calculate the number of pixels in the voting region | |
143 | - npts = nnz(M); | |
144 | - validPixels(v) = npts; | |
145 | - | |
146 | - %convert every index in the voting area from a local 3D index to a global 3D index (into the original image I) | |
147 | - global_px = vx + mx(pi); | |
148 | - global_py = vy + my(pi); | |
149 | - global_pz = vz + mz(pi); | |
150 | - | |
151 | - %convert the global 3D index of each point into a global 1D index | |
152 | - global_pi = sub2ind(ps, global_px, global_py, global_pz); | |
153 | - | |
154 | - g_v_prime (v, 1:npts) = global_pi; | |
155 | - | |
156 | - | |
157 | - Ivote( global_pi ) = Ivote( global_pi ) + vmag; | |
158 | - | |
159 | - end | |
160 | - | |
161 | - if itr ==1 | |
162 | - Ivote1 = single(Ivote); | |
163 | - | |
164 | - elseif itr ==2 | |
165 | - Ivote2 = single(Ivote); | |
166 | - | |
167 | - elseif itr ==3 | |
168 | - Ivote3 = single(Ivote); | |
169 | - | |
170 | - elseif itr ==4 | |
171 | - Ivote4 = single(Ivote); | |
172 | - | |
173 | - elseif itr == 5 | |
174 | - Ivote5 = single(Ivote); | |
175 | - end | |
176 | - t_v1 = toc; | |
177 | - disp(['voting done. time =',num2str(t_v1)]); | |
178 | - | |
179 | - % update the voting direction | |
180 | - if ang>0 | |
181 | - tic; | |
182 | - for v = 1: nV | |
183 | - % coordinates of the current voter | |
184 | - vx = Itx(v); | |
185 | - vy = Ity(v); | |
186 | - vz = Itz(v); | |
187 | - | |
188 | - %get the local value of the voting image | |
189 | - local_Ivote = Ivote(g_v_prime(v,1:validPixels(v))); | |
190 | - | |
191 | - %find the index of the maximum value | |
192 | - [~, local_max_idx] = max(local_Ivote); | |
193 | - | |
194 | - %convert this into a global subscript | |
195 | - [g_px, g_py, g_pz] = ind2sub(size(Ivote), g_v_prime(v,local_max_idx)); | |
196 | - | |
197 | - %compute the vector from the voter position to this position | |
198 | - Igrad_x(vx, vy, vz) = g_px - vx; | |
199 | - Igrad_y(vx, vy, vz) = g_py - vy; | |
200 | - Igrad_z(vx, vy, vz) = g_pz - vz; | |
201 | - | |
202 | - end | |
203 | - | |
204 | - | |
205 | - tdir1 = toc; | |
206 | - display (['updating dir done. time = ', num2str(tdir1)]); | |
207 | - ang = ang - d_ang; | |
208 | - end | |
209 | - end | |
210 | - | |
211 | - | |
212 | -%% | |
213 | -t = 350; | |
214 | -conn = [5 5 3]; | |
215 | -Icenter = local_max(Ivote, conn, t); | |
216 | -% center = Ivote1; | |
217 | -% center(center<t) = 0; | |
218 | -% center = imregionalmax(center); | |
219 | -% cn = nnz(center); | |
220 | -% [cx, cy, cz] = ind2sub(size(center), find(center)); | |
221 | -% Icenter = zeros(size(center)); | |
222 | -% for cc =1:cn | |
223 | -% Icenter(cx(cc), cy(cc), cz(cc)) = 255; | |
224 | -% end | |
225 | - | |
226 | -% fid_Ic = fopen('image_center2-300.vol', 'w'); | |
227 | -% fwrite(fid_Ic, Icenter); | |
228 | -% fclose(fid_Ic); | |
229 | -cn = nnz(Icenter); | |
230 | -[cx, cy, cz] = ind2sub(size(Icenter), find(Icenter)); | |
231 | -Ic2d = zeros(size(Icenter,1), size(Icenter,2)); | |
232 | -for cc =1:cn | |
233 | - Ic2d(cx(cc), cy(cc)) = 1; | |
234 | -end | |
235 | -I2d = max(I, [], 3); | |
236 | -% figure(1),imagesc(I2d); colormap(gray); | |
237 | -% figure(2),imagesc(Ic2d); colormap(gray); | |
238 | -% | |
239 | -out1(:,:,1) = mat2gray(I2d); | |
240 | -out1(:,:,2) = mat2gray(Ic2d); | |
241 | -out1(:,:,3) = mat2gray(I2d); | |
242 | -figure(1), imagesc((out1)); | |
243 | -%%% % imwrite(mat2gray(c2d), 'vote.bmp'); | |
244 | -%% | |
245 | -% figure(1); imagesc(squeeze(I(:,:,ceil(size(I,3)/2)))), colormap(gray); | |
246 | -% figure(33); imagesc(squeeze(Ivote3(:,:,ceil(size(Ivote,3)/2)))), colormap(gray); |
cpp/cudafunc.cu
... | ... | @@ -9,9 +9,11 @@ |
9 | 9 | #include "local_max3.cuh" |
10 | 10 | |
11 | 11 | |
12 | -void ivote3(float* center, float* img, float sigma[], float phi, float d_phi, unsigned int r[], | |
12 | +void ivote3(float* center, float* img, float sigma[], float anisotropy, float phi, float d_phi, unsigned int r[], | |
13 | 13 | int iter, float t, unsigned int conn[], unsigned int x, unsigned int y, unsigned int z){ |
14 | 14 | |
15 | + | |
16 | + cudaSetDevice(0); | |
15 | 17 | // compute the number of bytes in the input data |
16 | 18 | unsigned int bytes = x * y * z * sizeof(float); |
17 | 19 | |
... | ... | @@ -24,7 +26,7 @@ void ivote3(float* center, float* img, float sigma[], float phi, float d_phi, un |
24 | 26 | |
25 | 27 | //call the blurring function from the gpu. |
26 | 28 | gpu_gaussian_blur3<float>(gpuI0, sigma, x, y, z); |
27 | - | |
29 | + //cudaMemcpy(img, gpuI0, bytes, cudaMemcpyDeviceToHost); | |
28 | 30 | cudaDeviceSynchronize(); |
29 | 31 | |
30 | 32 | //assign memory on the gpu for the gradient along the X, y, z. |
... | ... | @@ -32,7 +34,7 @@ void ivote3(float* center, float* img, float sigma[], float phi, float d_phi, un |
32 | 34 | cudaMalloc(&gpu_grad, bytes*3); |
33 | 35 | |
34 | 36 | //call the gradient function from the gpu. |
35 | - gpu_gradient3<float>(gpu_grad, gpuI0, x, y, z); | |
37 | + gpu_gradient3<float>(gpu_grad, gpuI0, anisotropy, x, y, z); | |
36 | 38 | cudaFree(gpuI0); |
37 | 39 | |
38 | 40 | float* gpu_vote; |
... | ... | @@ -45,7 +47,7 @@ void ivote3(float* center, float* img, float sigma[], float phi, float d_phi, un |
45 | 47 | |
46 | 48 | gpu_vote3<float>(gpu_vote, gpu_grad, cos_phi, r, x, y, z); |
47 | 49 | cudaDeviceSynchronize(); |
48 | - if (i==0) | |
50 | + if (i==7) | |
49 | 51 | cudaMemcpy(img, gpu_vote, bytes, cudaMemcpyDeviceToHost); |
50 | 52 | |
51 | 53 | if (phi >= d_phi){ |
... | ... | @@ -58,7 +60,7 @@ void ivote3(float* center, float* img, float sigma[], float phi, float d_phi, un |
58 | 60 | } |
59 | 61 | |
60 | 62 | cudaFree(gpu_grad); |
61 | - cudaMemcpy(center, gpu_vote, bytes, cudaMemcpyDeviceToHost); | |
63 | + //cudaMemcpy(center, gpu_grad, bytes, cudaMemcpyDeviceToHost); | |
62 | 64 | |
63 | 65 | //allocate space on the gpu for the final detected cells. |
64 | 66 | float* gpu_output; |
... | ... | @@ -68,7 +70,7 @@ void ivote3(float* center, float* img, float sigma[], float phi, float d_phi, un |
68 | 70 | gpu_local_max3<float>(gpu_output, gpu_vote, t, conn, x, y, z); |
69 | 71 | |
70 | 72 | //copy the final result to the cpu. |
71 | - //cudaMemcpy(center, gpu_output, bytes, cudaMemcpyDeviceToHost); | |
73 | + cudaMemcpy(center, gpu_output, bytes, cudaMemcpyDeviceToHost); | |
72 | 74 | |
73 | 75 | |
74 | 76 | cudaFree(gpu_vote); | ... | ... |
cpp/gradient3.cuh
... | ... | @@ -7,7 +7,7 @@ |
7 | 7 | #include <stim/cuda/cudatools/error.h> |
8 | 8 | |
9 | 9 | template<typename T> |
10 | -__global__ void gradient3(T* out, T* in, int x, int y, int z){ | |
10 | +__global__ void gradient3(T* out, T* in, float anisotropy, int x, int y, int z){ | |
11 | 11 | |
12 | 12 | //calculate x,y,z coordinates for this thread |
13 | 13 | int xi = blockIdx.x * blockDim.x + threadIdx.x; |
... | ... | @@ -55,11 +55,13 @@ __global__ void gradient3(T* out, T* in, int x, int y, int z){ |
55 | 55 | if(zi > 0 && zi < z-1) |
56 | 56 | out[i * 3 + 2] = (in[i_zp] - in[i_zn]) / 2; |
57 | 57 | |
58 | + out[i * 3 + 2] *= 1/anisotropy; | |
59 | + | |
58 | 60 | } |
59 | 61 | |
60 | 62 | template<typename T> |
61 | 63 | |
62 | -void gpu_gradient3(T* gpuGrad, T* gpuI, unsigned int x, unsigned int y, unsigned int z){ | |
64 | +void gpu_gradient3(T* gpuGrad, T* gpuI, float anisotropy, unsigned int x, unsigned int y, unsigned int z){ | |
63 | 65 | |
64 | 66 | |
65 | 67 | int max_threads = stim::maxThreadsPerBlock(); |
... | ... | @@ -67,12 +69,12 @@ void gpu_gradient3(T* gpuGrad, T* gpuI, unsigned int x, unsigned int y, unsigned |
67 | 69 | dim3 blocks(x / threads.x + 1, (y / threads.y + 1) * z); |
68 | 70 | |
69 | 71 | //call the GPU kernel to determine the gradient |
70 | - gradient3<T> <<< blocks, threads >>>(gpuGrad, gpuI, x, y, z); | |
72 | + gradient3<T> <<< blocks, threads >>>(gpuGrad, gpuI, anisotropy, x, y, z); | |
71 | 73 | |
72 | 74 | } |
73 | 75 | |
74 | 76 | template<typename T> |
75 | -void cpu_gradient3(T* out, T* in, unsigned int x, unsigned int y, unsigned int z){ | |
77 | +void cpu_gradient3(T* out, T* in, float anisotropy, unsigned int x, unsigned int y, unsigned int z){ | |
76 | 78 | |
77 | 79 | //get the number of pixels in the image |
78 | 80 | unsigned int pixels = x * y * z; |
... | ... | @@ -90,7 +92,7 @@ void cpu_gradient3(T* out, T* in, unsigned int x, unsigned int y, unsigned int z |
90 | 92 | cudaMalloc(&gpuOut, bytes * 3); //the output image will have two channels (x, y) |
91 | 93 | |
92 | 94 | //call the GPU version of this function |
93 | - gpu_gradient3(gpuOut, gpuIn, x, y, z); | |
95 | + gpu_gradient3(gpuOut, gpuIn, anisotropy, x, y, z); | |
94 | 96 | |
95 | 97 | //copy the results to the CPU |
96 | 98 | cudaMemcpy(out, gpuOut, bytes * 3, cudaMemcpyDeviceToHost); | ... | ... |
cpp/local_max3.cuh
... | ... | @@ -36,7 +36,7 @@ __global__ void cuda_local_max3(T* gpu_center, T* gpu_vote, T t, int conn_x, int |
36 | 36 | |
37 | 37 | if (xl>=0 && yl>=0 && zl>=0 && xl<x && yl<y && zl<z){ |
38 | 38 | |
39 | - unsigned int i_l = zl * x * y + yl * x + xl; | |
39 | + int i_l = zl * x * y + yl * x + xl; | |
40 | 40 | if (gpu_vote[i_l] > lv_i) return; |
41 | 41 | } |
42 | 42 | } | ... | ... |
cpp/main.cpp
... | ... | @@ -11,7 +11,7 @@ |
11 | 11 | #define pi 3.14159 |
12 | 12 | |
13 | 13 | |
14 | -void ivote3(float* center, float* img, float std[], float phi, float d_phi, unsigned int r[], int iter, float t, unsigned int conn[], | |
14 | +void ivote3(float* center, float* img, float std[], float anisotropy, float phi, float d_phi, unsigned int r[], int iter, float t, unsigned int conn[], | |
15 | 15 | unsigned int x, unsigned int y, unsigned int z); |
16 | 16 | |
17 | 17 | void invert_data(float* cpuI, unsigned int x, unsigned int y, unsigned int z){ |
... | ... | @@ -24,9 +24,23 @@ void invert_data(float* cpuI, unsigned int x, unsigned int y, unsigned int z){ |
24 | 24 | } |
25 | 25 | } |
26 | 26 | } |
27 | - | |
27 | + | |
28 | 28 | int main(int argc, char** argv){ |
29 | 29 | |
30 | + | |
31 | + cudaDeviceProp prop; | |
32 | + int count; | |
33 | + cudaGetDeviceCount(&count); | |
34 | + //printf("cudadevicecount: %i\n", count); | |
35 | + for (int i=0; i<count; i++){ | |
36 | + cudaGetDeviceProperties(&prop, i); | |
37 | + printf("current device ID: %d\n", i); | |
38 | + printf("device name: %s\n", prop.name); | |
39 | + printf("total global mem: %lu\n", prop.totalGlobalMem); | |
40 | + } | |
41 | + | |
42 | + | |
43 | + | |
30 | 44 | //output advertisement |
31 | 45 | std::cout<<std::endl<<std::endl; |
32 | 46 | std::cout<<"========================================================================="<<std::endl; |
... | ... | @@ -48,6 +62,8 @@ int main(int argc, char** argv){ |
48 | 62 | args.add("y", "size of the dataset along Y axis", "positive value"); |
49 | 63 | args.add("z", "size of the dataset along Z axis", "positive value"); |
50 | 64 | args.add("t", "threshold value for the final result", "positive valu"); |
65 | + args.add("invert", "to invert the input data set", "string"); | |
66 | + args.add("anisotropy", "anisotropy value of the imaging", "positive value"); | |
51 | 67 | //parse the command line arguments. |
52 | 68 | args.parse(argc, argv); |
53 | 69 | |
... | ... | @@ -81,14 +97,18 @@ int main(int argc, char** argv){ |
81 | 97 | |
82 | 98 | //set the threshold. |
83 | 99 | float t = args["t"].as_float(); |
84 | - | |
85 | - unsigned int r[3] = { 9, 9, 5}; | |
86 | - float sigma[3] = { 3, 3, 1.5}; | |
87 | - unsigned int conn[3] = { 5, 5, 3}; | |
88 | - float phi_deg = 20.1; | |
100 | + //set the anisotropy | |
101 | + float anisotropy = args["anisotropy"].as_float(); | |
102 | + unsigned int rmax = 10 ; | |
103 | + unsigned int r[3] = { rmax, rmax, rmax}; | |
104 | + float std = 5; | |
105 | + float sigma[3] = { std, std, std}; | |
106 | + unsigned int nlmax = 5; | |
107 | + unsigned int conn[3] = { nlmax, nlmax, nlmax}; | |
108 | + float phi_deg = 25.0; | |
89 | 109 | float phi = phi_deg * pi /180; |
90 | - int iter = 2; | |
91 | - float d_phi = phi/(iter -1); | |
110 | + int iter = 8; | |
111 | + float d_phi = phi/(iter+2); | |
92 | 112 | |
93 | 113 | std::string filename = Ifilename.str(); |
94 | 114 | unsigned int bytes = x*y*z*sizeof(float); |
... | ... | @@ -100,31 +120,56 @@ int main(int argc, char** argv){ |
100 | 120 | std::ifstream nissl(filename, std::ios::in | std::ios::binary); |
101 | 121 | nissl.read((char*)cpuI, bytes); |
102 | 122 | nissl.close(); |
103 | - | |
104 | - invert_data(cpuI, x, y, z); | |
123 | + if(args["invert"].is_set()) | |
124 | + invert_data(cpuI, x, y, z); | |
105 | 125 | |
106 | 126 | //write a new file from the cpuI. |
107 | - std::ofstream original("output/original_invert--512.vol", std::ofstream::out | std::ofstream::binary); | |
127 | + std::ofstream original("std5.5-r10.10-v8/inv-128.vol", std::ofstream::out | std::ofstream::binary); | |
108 | 128 | original.write((char*)cpuI, bytes); |
109 | 129 | original.close(); |
110 | 130 | |
111 | 131 | //allocate space on the cpu for the output result |
112 | - float* cpu_out = (float*) malloc(bytes); | |
132 | + float* cpu_out = (float*) malloc(bytes*3); | |
113 | 133 | |
114 | 134 | // call the ivote function |
115 | - ivote3(cpu_out, cpuI, sigma, phi, d_phi, r, iter, t, conn, x, y, z); | |
135 | + ivote3(cpu_out, cpuI, sigma, anisotropy, phi, d_phi, r, iter, t, conn, x, y, z); | |
116 | 136 | |
117 | 137 | //write the blurred file from the cpuI. |
118 | - std::ofstream fblur("output/v1--512.vol", std::ofstream::out | std::ofstream::binary); | |
138 | + std::ofstream fblur("vote-check/vote8.vol", std::ofstream::out | std::ofstream::binary); | |
119 | 139 | fblur.write((char*)cpuI, bytes); |
120 | 140 | fblur.close(); |
121 | - | |
141 | + /* | |
142 | + stim::image<float>imgrad3; | |
143 | + imgrad3.set_interleaved3(cpu_out, 128,128,128,3); | |
144 | + std::ofstream fgx("syn/gx-128.vol", std::ofstream::out | std::ofstream::binary); | |
145 | + fgx.write((char*)imgrad3.channel(0).data(), bytes); | |
146 | + fgx.close(); | |
147 | + */ | |
122 | 148 | //write the output file. |
123 | - std::ofstream fo("output/" + OutName.str(), std::ofstream::out | std::ofstream::binary); | |
149 | + std::ofstream fo("std5.5-r10.10-v8/" + OutName.str(), std::ofstream::out | std::ofstream::binary); | |
124 | 150 | fo.write((char*)cpu_out, bytes); |
125 | 151 | fo.close(); |
126 | 152 | |
127 | - | |
153 | + // creat a file for saving the list centers | |
154 | + | |
155 | + std::ofstream list("std5.5-r10.10-v8/center.txt"); | |
156 | + if (list.is_open()){ | |
157 | + | |
158 | + for (int ix=0; ix<x; ix++){ | |
159 | + for (int iy=0; iy<y; iy++){ | |
160 | + for (int iz=0; iz<z; iz++){ | |
161 | + | |
162 | + int idx = iz * x * y + iy * x + ix; | |
163 | + if (cpu_out[idx]==1){ | |
164 | + list << ix << "\t" << iy << "\t"<< iz << '\n' ; | |
165 | + | |
166 | + } | |
167 | + } | |
168 | + } | |
169 | + } | |
170 | + | |
171 | + list.close(); | |
172 | + } | |
128 | 173 | |
129 | 174 | cudaDeviceReset(); |
130 | 175 | ... | ... |
cpp/update_dir3.cuh
... | ... | @@ -32,7 +32,7 @@ |
32 | 32 | // define a local variable to maximum value of the vote image in the voting area for this voter |
33 | 33 | float max = 0; |
34 | 34 | float l_vote = 0; |
35 | - // define local variables for the x, y, and z coordinations where the maximum happened | |
35 | + // define local variables for the x, y, and z coordinations point to the vote direction | |
36 | 36 | float id_x = g_v_x; |
37 | 37 | float id_y = g_v_y; |
38 | 38 | float id_z = g_v_z; | ... | ... |
cpp/vote3.cuh
... | ... | @@ -51,6 +51,7 @@ |
51 | 51 | float g_v_y = gpu_grad[id_v * 3 + 1]; |
52 | 52 | float g_v_z = gpu_grad[id_v * 3 + 2]; |
53 | 53 | float mag_v = sqrt( g_v_x * g_v_x + g_v_y * g_v_y + g_v_z * g_v_z); |
54 | + | |
54 | 55 | //calculate the distance between the pixel and the current voter. |
55 | 56 | float x_sq = x_v * x_v; |
56 | 57 | float y_sq = y_v * y_v; | ... | ... |