1 |
__kernel void raycast (__global float * img, |
__kernel void raycast ( float * img, |
2 |
__global float * h1, |
float * h1, |
3 |
__global float * h2, |
float * h2, |
4 |
__global float * out, |
float * out, |
5 |
float4 orig, |
float4 orig, |
6 |
float4 eyeVector, |
float4 eyeVector, |
7 |
float4 cVec, |
float4 cVec, |
9 |
float16 transformMatrix, |
float16 transformMatrix, |
10 |
float stepSize, |
float stepSize, |
11 |
int length, |
int length, |
12 |
int width) |
int width, |
13 |
|
int height) |
14 |
{ |
{ |
15 |
int row = get_global_id(0), col = get_global_id(1); |
int row = get_global_id(0), col = get_global_id(1); |
16 |
|
|
22 |
int4 n1; |
int4 n1; |
23 |
int4 n; |
int4 n; |
24 |
|
|
|
|
|
|
|
|
25 |
float4 d = (float4) (h2[0],h1[0],h1[0],h2[0]); |
float4 d = (float4) (h2[0],h1[0],h1[0],h2[0]); |
26 |
float4 c = (float4) (h2[1],h1[1],h1[1],h2[1]); |
float4 c = (float4) (h2[1],h1[1],h1[1],h2[1]); |
27 |
float4 b = (float4) (h2[2],h1[2],h1[2],h2[2]); |
float4 b = (float4) (h2[2],h1[2],h1[2],h2[2]); |
29 |
|
|
30 |
float4 pos = orig + (float)row * rVec + (float)col * cVec; |
float4 pos = orig + (float)row * rVec + (float)col * cVec; |
31 |
float4 dir = (pos - eyeVector) / fabs(pos - eyeVector); |
float4 dir = (pos - eyeVector) / fabs(pos - eyeVector); |
32 |
|
float4 NSize = (float4) ((float)height,(float)length,(float)width,1.0f); |
33 |
|
|
34 |
pt.w = 1.0f; |
pt.w = 1.0f; |
35 |
pos.w = 1.0f; |
pos.w = 1.0f; |
50 |
dot(pos,transformMatrix.scdef)); |
dot(pos,transformMatrix.scdef)); |
51 |
|
|
52 |
|
|
|
|
|
53 |
f.xyzw = (float4) (modf(imgPt.x,&x),modf(imgPt.y,&y),modf(imgPt.z,&z),1.0f); |
f.xyzw = (float4) (modf(imgPt.x,&x),modf(imgPt.y,&y),modf(imgPt.z,&z),1.0f); |
54 |
|
|
55 |
n1 = (int4) ((int)x,(int)y,(int)z,1.0f); |
n1 = (int4) ((int)x,(int)y,(int)z,1.0f); |
56 |
|
|
57 |
if( (n1.x > 1 && n1.x < 11) && |
if( (imgPt.x > 1 && imgPt.x < NSize.x - 2) && |
58 |
(n1.y > 1 && n1.y < 12) && |
(imgPt.y > 1 && imgPt.y < NSize.y - 2) && |
59 |
(n1.z > 1 && n1.z < 14)) |
(imgPt.z > 1 && imgPt.z < NSize.z - 2)) |
60 |
{ |
{ |
61 |
|
|
62 |
|
|
63 |
|
|
64 |
// the t value for h(fx - i) |
// the t value for h(fx - i) |
65 |
t_i = (float4) (-1.0-f.x, -f.x,f.x - 1.0, f.x - 2.0); |
t_i = (float4) (-1.0-f.x, -f.x,f.x - 1.0, f.x - 2.0); |
66 |
|
|
82 |
((n1.x+2) * length * width + n1.y+2 * width + n1.z+2)); |
((n1.x+2) * length * width + n1.y+2 * width + n1.z+2)); |
83 |
|
|
84 |
|
|
85 |
|
|
86 |
v = (float4)(img[n.x], |
v = (float4)(img[n.x], |
87 |
img[n.y], |
img[n.y], |
88 |
img[n.z], |
img[n.z], |
91 |
|
|
92 |
probedVal = dot(v,value); // V(n + <i,j,k>) * summations of the h(x) components |
probedVal = dot(v,value); // V(n + <i,j,k>) * summations of the h(x) components |
93 |
|
|
|
// End Probe Operation |
|
94 |
|
|
95 |
|
// End Probe Operation |
96 |
if(maxValue < probedVal) |
if(maxValue < probedVal) |
97 |
maxValue = probedVal; |
maxValue = probedVal; |
98 |
} |
} |
|
if(row == 0 && col == 0) |
|
|
printf("IMG PT (%f,%f,%f)\n", imgPt.x,imgPt.y,imgPt.z ); |
|
99 |
} |
} |
100 |
|
if(row == 0 && col == 74) |
101 |
|
printf("Max Value: %f\n", maxValue ); |
102 |
out[row * 200 + col] = maxValue; |
out[row * 200 + col] = maxValue; |
103 |
} |
} |
104 |
} |
} |