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

Parent Directory | Revision Log

---

Revision 3349 - (view) (download)

1 :	lamonts	302	/* mip.cl
2 :			*
3 :	jhr	3349	* This code is part of the Diderot Project (http://diderot-language.cs.uchicago.edu)
4 :			*
5 :			* COPYRIGHT (c) 2015 The University of Chicago
6 :	lamonts	302	* All rights reserved.
7 :			*
8 :			* An OpenCL file containing a probe kernel and mip kernel implementation
9 :			*/
10 :
11 :			__constant float h[4][4] = { // bspln3
12 :			{ 1.33333, 2.0, 1.0, 0.166667 }, // -2 .. -1
13 :			{ 0.666667, 0.0, -1.0, -0.5 }, // -1 .. 0
14 :			{ 0.666667, 0.0, -1.0, 0.5 }, // 0 .. 1
15 :			{ 1.33333, -2.0, 1.0, -0.166667 }, // 1 .. 2
16 :			};
17 :
18 :	jhr	310	//__constant float4 eyeVec = (float4)(25,15,10,1);
19 :			//__constant float4 origVec = (float4)(8.83877,2.5911,7.65275,0);
20 :			//__constant float4 cVec= (float4)(-0.0151831,0.0278357,0,0);
21 :			//__constant float4 rVec = (float4)(0.0074887,0.00408474,-0.0305383,0);
22 :			__constant float4 eyeVec = (float4)(127.331, -1322.05, 272.53, 0);
23 :			__constant float4 origVec = (float4)(122.835,17.7112,188.044, 0);
24 :			__constant float4 cVec= (float4)(-0.00403611,-0.029826,-0.244066, 0);
25 :			__constant float4 rVec = (float4)(-0.245595,-0.0112916,0.00544129, 0);
26 :			__constant float stepSize = 0.5;
27 :	lamonts	302	__constant int s = 2;
28 :
29 :
30 :	jhr	441	__kernel float probe(float * img, int * sAxis, float4 imgPos)
31 :			{
32 :	lamonts	302
33 :	jhr	441	float probedVal;
34 :			float4 f, nf, t, hx, hy, hz;
35 :			int4 n;
36 :	lamonts	302
37 :	jhr	441	float4 d = (float4) (h[3][0],h[2][0],h[1][0],h[0][0]);
38 :			float4 c = (float4) (h[3][1],h[2][1],h[1][1],h[0][1]);
39 :			float4 b = (float4) (h[3][2],h[2][2],h[1][2],h[0][2]);
40 :			float4 a = (float4) (h[3][3],h[2][3],h[1][3],h[0][3]);
41 :	lamonts	302
42 :	jhr	441	f = modf(imgPos,&nf);
43 :			n = convert_int4(nf);
44 :
45 :	lamonts	302
46 :	jhr	441	t = (float4) (f.x + 1, f.x, f.x -1, f.x - 2);
47 :			hx = d + t * (c + t * (b + t*a));
48 :			t = (float4) (f.y + 1, f.y, f.y - 1, f.y - 2);
49 :			hy = d + t * (c + t * (b + t*a));
50 :			t = (float4) (f.z + 1, f.z, f.z - 1, f.z -2);
51 :			hz = d + t * (c + t * (b + t*a));
52 :	lamonts	302
53 :	jhr	441	float vx[4];
54 :			float vy[4];
55 :
56 :			for(int k = 1-s; k <= s; k++) { // z is the slowest dimension
57 :			for(int j = 1-s; j <= s; j++) { // y is the medium dimension
58 :			int index = sAxis[0]*sAxis[1]*(n.z+k) + sAxis[0]*(n.y+j) + (n.x-1);
59 :			float4 v = (float4)(img[index], img[index+1], img[index+2], img[index+3]);
60 :			vx[j+s-1] = dot(v,hx);
61 :			}
62 :			vy[k+s-1] = dot ((float4) (vx[0],vx[1],vx[2],vx[3]),hy);
63 :			}
64 :	lamonts	302
65 :	jhr	441	probedVal = dot((float4) (vy[0],vy[1],vy[2],vy[3]),hz);
66 :	lamonts	302
67 :	jhr	441	return probedVal;
68 :	lamonts	302
69 :			}
70 :	jhr	441
71 :	lamonts	302	__kernel void mip ( float * img,
72 :	jhr	441	float * out,
73 :			float16 transformMatrix,
74 :			int2 workDim,
75 :			int * sAxis)
76 :			{
77 :	lamonts	302
78 :	jhr	441	int row = get_global_id(0), col = get_global_id(1);
79 :	lamonts	302
80 :	jhr	441	if(row <= workDim.x && col <= workDim.y) {
81 :			float t, probedVal, maxValue = -INFINITY;
82 :			float4 imgPt;
83 :	lamonts	177
84 :	jhr	441	float4 pos = origVec + (float)row * rVec + (float)col * cVec;
85 :			float4 dir = normalize(pos - eyeVec);
86 :	lamonts	177
87 :	jhr	441	pos.w = 0.0;
88 :			dir.w = 0.0;
89 :	lamonts	177
90 :	jhr	441	for(t = 0.0; t < 200; t+= stepSize) {
91 :			pos = pos + stepSize * dir;
92 :
93 :			// Transform the value to image space position.
94 :			imgPt = (float4) (dot(pos,transformMatrix.s0123),
95 :			dot(pos,transformMatrix.s4567),
96 :			dot(pos,transformMatrix.s89ab),
97 :			dot(pos,transformMatrix.scdef));
98 :
99 :			if ((imgPt.x > 1 && imgPt.x < (sAxis[0] - 2)
100 :			&& (imgPt.y > 1 && imgPt.y < (sAxis[1] - 2))
101 :			&& (imgPt.z > 1 && imgPt.z < (sAxis[2] - 2))) {
102 :			probedVal = probe(img,sAxis,imgPt);
103 :			if (maxValue < probedVal)
104 :			maxValue = probedVal;
105 :			}
106 :			}
107 :			out[row * workDim.x + col] = maxValue;
108 :			}
109 :	jhr	435	}

---

root@smlnj-gforge.cs.uchicago.edu	ViewVC Help
Powered by ViewVC 1.0.0