SCM Repository
Annotation of /trunk/test/probe/probe.cl
Parent Directory
|
Revision Log
Revision 310 - (view) (download)
1 : | jhr | 310 | __kernel void diderot (float * img, |
2 : | jhr | 262 | float * h0, |
3 : | float * h1, | ||
4 : | float * h2, | ||
5 : | float * h3, | ||
6 : | float * posValues, | ||
7 : | float * out, | ||
8 : | float16 transformMatrix, | ||
9 : | int ySize, | ||
10 : | int zSize, | ||
11 : | int xSize) | ||
12 : | lamonts | 259 | { |
13 : | lamonts | 277 | int location = get_global_id(0); |
14 : | float4 pos = (float4)(posValues[location *3], | ||
15 : | jhr | 262 | posValues[(location *3) + 1], |
16 : | posValues[(location *3) + 2],1.0f); | ||
17 : | |||
18 : | float probedVal; | ||
19 : | float4 f, imgPt,nf,v, hx, hy, hz, h_tx, h_ty, hMult; | ||
20 : | int4 n1; | ||
21 : | lamonts | 277 | |
22 : | float4 d = (float4) (h3[0],h2[0],h1[0],h0[0]); | ||
23 : | float4 c = (float4) (h3[1],h2[1],h1[1],h0[1]); | ||
24 : | float4 b = (float4) (h3[2],h2[2],h1[2],h0[2]); | ||
25 : | float4 a = (float4) (h3[3],h2[3],h1[3],h0[3]); | ||
26 : | jhr | 262 | |
27 : | // Transform the value to image space position. | ||
28 : | imgPt = (float4) (dot(pos,transformMatrix.s0123), | ||
29 : | dot(pos,transformMatrix.s4567), | ||
30 : | dot(pos,transformMatrix.s89ab), | ||
31 : | dot(pos,transformMatrix.scdef)); | ||
32 : | |||
33 : | |||
34 : | nf = floor(imgPt); | ||
35 : | f = imgPt - nf; | ||
36 : | n1 = convert_int4(nf); | ||
37 : | lamonts | 277 | |
38 : | float4 t; | ||
39 : | int s = 2; | ||
40 : | |||
41 : | t = (float4) (f.x + 1, f.x, f.x -1, f.x - 2); | ||
42 : | hx = a + t * (b + t * (c + t*d)); | ||
43 : | t = (float4) (f.y + 1, f.y, f.y - 1, f.y - 2); | ||
44 : | hy = a + t * (b + t * (c + t*d)); | ||
45 : | t = (float4) (f.z + 1, f.z, f.z - 1, f.z -2); | ||
46 : | hz = a + t * (b + t * (c + t*d)); | ||
47 : | |||
48 : | float vy[4], vz[4]; | ||
49 : | |||
50 : | for(int i = 1-s; i <= s; i++) { | ||
51 : | for(int j = 1-s; j <= s; j++) { | ||
52 : | float4 v = (float4) ( | ||
53 : | img[ (n1.x+i) * ySize * zSize + (n1.y+j) * zSize + n1.z-1], | ||
54 : | img[ (n1.x+i) * ySize * zSize + (n1.y+j) * zSize + n1.z], | ||
55 : | img[ (n1.x+i) * ySize * zSize + (n1.y+j) * zSize + n1.z+1], | ||
56 : | img[ (n1.x+i) * ySize * zSize + (n1.y+j) * zSize + n1.z+2]); | ||
57 : | vz[j+s-1] += dot(v,hz); | ||
58 : | } | ||
59 : | vy[i+s-1] = dot ((float4) (vz[0],vz[1],vz[2],vz[3]),hy); | ||
60 : | } | ||
61 : | |||
62 : | probedVal = dot((float4) (vy[0],vy[1],vy[2],vy[3]),hx); | ||
63 : | |||
64 : | printf("Probed @Pos(%.2f,%.2f,%.2f): %f\n",pos.x,pos.y,pos.z,probedVal); | ||
65 : | |||
66 : | out[location] = probedVal; | ||
67 : | |||
68 : | /*** ===== Previous Verison ====== | ||
69 : | jhr | 262 | |
70 : | // the t value for h(fx - i) | ||
71 : | float4 t_i = (float4) (f.x + 1.0, f.x, f.x - 1.0, f.x - 2.0); | ||
72 : | |||
73 : | // the t value for h(fy - j) | ||
74 : | float4 t_j = (float4) (f.y + 1.0, f.y, f.y - 1.0, f.y - 2.0); | ||
75 : | |||
76 : | float4 t_k = (float4) (f.z + 1.0, f.z, f.z - 1.0, f.z - 2.0); | ||
77 : | float4 h_k = d + t_k * (c + t_k * (b + t_k * a)); | ||
78 : | lamonts | 277 | float tx[4], ty[4]; |
79 : | |||
80 : | for (int i = -1; i <= 2; i++) { | ||
81 : | for (int j = -1; j <= 2; j++) { | ||
82 : | // compute z axis using vectors | ||
83 : | int index = (n1.x+i) * ySize * zSize + (n1.y+j) * zSize + n1.z-1; | ||
84 : | v = (float4)( | ||
85 : | img[index], | ||
86 : | img[index+1], | ||
87 : | img[index+2], | ||
88 : | img[index+3]); | ||
89 : | ty[j+1] = dot(v, h_k); | ||
90 : | } | ||
91 : | tx[i+1] = dot((float4)(ty[0], ty[1], ty[2], ty[3]), d + t_j * (c + t_j * (b + t_j * a))); | ||
92 : | } | ||
93 : | probedVal = dot((float4)(tx[0],tx[1],tx[2],tx[3]), d + t_i * (c + t_i * (b + t_i * a))); | ||
94 : | jhr | 262 | |
95 : | lamonts | 277 | === End of Previous Verison === */ |
96 : | |||
97 : | jhr | 310 | } |
root@smlnj-gforge.cs.uchicago.edu | ViewVC Help |
Powered by ViewVC 1.0.0 |