SCM Repository
Annotation of /trunk/test/MIP/mip.cl
Parent Directory
|
Revision Log
Revision 302 - (view) (download)
1 : | lamonts | 302 | /* mip.cl |
2 : | * | ||
3 : | * COPYRIGHT (c) 2010 The Diderot Project (http://diderot.cs.uchicago.edu) | ||
4 : | * All rights reserved. | ||
5 : | * | ||
6 : | * An OpenCL file containing a probe kernel and mip kernel implementation | ||
7 : | */ | ||
8 : | |||
9 : | __constant float h[4][4] = { // bspln3 | ||
10 : | { 1.33333, 2.0, 1.0, 0.166667 }, // -2 .. -1 | ||
11 : | { 0.666667, 0.0, -1.0, -0.5 }, // -1 .. 0 | ||
12 : | { 0.666667, 0.0, -1.0, 0.5 }, // 0 .. 1 | ||
13 : | { 1.33333, -2.0, 1.0, -0.166667 }, // 1 .. 2 | ||
14 : | }; | ||
15 : | |||
16 : | __constant float4 eyeVec = (float4)(25,15,10,1); | ||
17 : | __constant float4 origVec = (float4)(8.83877,2.5911,7.65275,0); | ||
18 : | __constant float4 cVec= (float4)(-0.0151831,0.0278357,0,0); | ||
19 : | __constant float4 rVec = (float4)(0.0074887,0.00408474,-0.0305383,0); | ||
20 : | __constant float stepSize = 0.1; | ||
21 : | __constant int s = 2; | ||
22 : | |||
23 : | |||
24 : | __kernel float probe(float * img, int * sAxis, float4 imgPos) { | ||
25 : | |||
26 : | float probedVal; | ||
27 : | float4 f, nf, t, hx, hy, hz; | ||
28 : | int4 n; | ||
29 : | |||
30 : | float4 d = (float4) (h[3][0],h[2][0],h[1][0],h[0][0]); | ||
31 : | float4 c = (float4) (h[3][1],h[2][1],h[1][1],h[0][1]); | ||
32 : | float4 b = (float4) (h[3][2],h[2][2],h[1][2],h[0][2]); | ||
33 : | float4 a = (float4) (h[3][3],h[2][3],h[1][3],h[0][3]); | ||
34 : | |||
35 : | f = modf(imgPos,&nf); | ||
36 : | n = convert_int4(nf); | ||
37 : | |||
38 : | |||
39 : | t = (float4) (f.x + 1, f.x, f.x -1, f.x - 2); | ||
40 : | hx = d + t * (c + t * (b + t*a)); | ||
41 : | t = (float4) (f.y + 1, f.y, f.y - 1, f.y - 2); | ||
42 : | hy = d + t * (c + t * (b + t*a)); | ||
43 : | t = (float4) (f.z + 1, f.z, f.z - 1, f.z -2); | ||
44 : | hz = d + t * (c + t * (b + t*a)); | ||
45 : | |||
46 : | float vy[4] = {0.0f,0.0f,0.0f,0.0f}; | ||
47 : | float vz[4] = {0.0f,0.0f,0.0f,0.0f}; | ||
48 : | |||
49 : | for(int i = 1-s; i <= s; i++) { | ||
50 : | for(int j = 1-s; j <= s; j++) { | ||
51 : | float4 v = (float4) ( | ||
52 : | img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z-1))], | ||
53 : | img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z))], | ||
54 : | img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z+1))], | ||
55 : | img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z+2))]); | ||
56 : | vz[j+s-1] += dot(v,hz); | ||
57 : | } | ||
58 : | vy[i+s-1] = dot ((float4) (vz[0],vz[1],vz[2],vz[3]),hy); | ||
59 : | } | ||
60 : | |||
61 : | probedVal = dot((float4) (vy[0],vy[1],vy[2],vy[3]),hx); | ||
62 : | |||
63 : | return probedVal; | ||
64 : | |||
65 : | } | ||
66 : | __kernel void mip ( float * img, | ||
67 : | float * out, | ||
68 : | float16 transformMatrix, | ||
69 : | int2 workDim, | ||
70 : | int * sAxis){ | ||
71 : | |||
72 : | lamonts | 177 | int row = get_global_id(0), col = get_global_id(1); |
73 : | lamonts | 302 | |
74 : | if(row <= workDim.x && col <= workDim.y) | ||
75 : | { | ||
76 : | float t, probedVal, maxValue = -INFINITY; | ||
77 : | float4 imgPt; | ||
78 : | lamonts | 177 | |
79 : | lamonts | 302 | float4 pos = origVec + (float)row * rVec + (float)col * cVec; |
80 : | float4 dir = normalize(pos - eyeVec); | ||
81 : | jhr | 218 | |
82 : | lamonts | 302 | pos.w = 0.0; |
83 : | dir.w = 0.0; | ||
84 : | |||
85 : | jhr | 218 | for(t = 0.0; t < 40; t+= stepSize) |
86 : | lamonts | 302 | { |
87 : | lamonts | 177 | pos = pos + stepSize * dir; |
88 : | |||
89 : | // Transform the value to image space position. | ||
90 : | imgPt = (float4) (dot(pos,transformMatrix.s0123), | ||
91 : | jhr | 218 | dot(pos,transformMatrix.s4567), |
92 : | dot(pos,transformMatrix.s89ab), | ||
93 : | dot(pos,transformMatrix.scdef)); | ||
94 : | lamonts | 177 | |
95 : | lamonts | 302 | |
96 : | if( (imgPt.x > 1 && imgPt.x < (sAxis[0] - 2)) && | ||
97 : | (imgPt.y > 1 && imgPt.y < (sAxis[1]- 2)) && | ||
98 : | (imgPt.z > 1 && imgPt.z < (sAxis[2] - 2))) | ||
99 : | lamonts | 191 | { |
100 : | lamonts | 302 | probedVal = probe(img,sAxis,imgPt); |
101 : | jhr | 218 | if(maxValue < probedVal) |
102 : | lamonts | 203 | maxValue = probedVal; |
103 : | lamonts | 191 | } |
104 : | lamonts | 203 | } |
105 : | lamonts | 302 | out[row * workDim.x + col] = maxValue; |
106 : | lamonts | 177 | } |
107 : | } |
root@smlnj-gforge.cs.uchicago.edu | ViewVC Help |
Powered by ViewVC 1.0.0 |