Home My Page Projects Code Snippets Project Openings diderot
Summary Activity Tracker Tasks SCM

SCM Repository

[diderot] Annotation of /trunk/test/MIP/mip.cl
ViewVC logotype

Annotation of /trunk/test/MIP/mip.cl

Parent Directory Parent Directory | Revision Log 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