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

SCM Repository

[diderot] Diff of /branches/pure-cfg/test/MIP/mip.cl
ViewVC logotype

Diff of /branches/pure-cfg/test/MIP/mip.cl

Parent Directory Parent Directory | Revision Log Revision Log | View Patch Patch

revision 203, Tue Aug 3 14:46:27 2010 UTC revision 441, Wed Oct 20 19:21:32 2010 UTC
# Line 1  Line 1 
1  __kernel void raycast ( float * img,  /* mip.cl
2                                              float * h1,   *
3                                              float * h2,   * COPYRIGHT (c) 2010 The Diderot Project (http://diderot-language.cs.uchicago.edu)
4                                              float * out,   * All rights reserved.
5                                              float4  orig,   *
6                                              float4  eyeVector,   * An OpenCL file containing a probe kernel and mip kernel implementation
7                                              float4  cVec,   */
8                                              float4  rVec,  
9                                              float16  transformMatrix,  __constant float h[4][4] = {            // bspln3
10                                              float stepSize,          { 1.33333,   2.0,  1.0,  0.166667 },   // -2 .. -1
11                                              int length,          { 0.666667,  0.0, -1.0, -0.5 },        // -1 .. 0
12                                              int width,          { 0.666667,  0.0, -1.0,  0.5 },        //  0 .. 1
13                                              int height)          { 1.33333,  -2.0,  1.0, -0.166667 },   //  1 .. 2
14  {  };
15          int row = get_global_id(0), col = get_global_id(1);  
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 float4 eyeVec = (float4)(127.331, -1322.05, 272.53, 0);
21    __constant float4 origVec = (float4)(122.835,17.7112,188.044, 0);
22    __constant float4 cVec= (float4)(-0.00403611,-0.029826,-0.244066, 0);
23    __constant float4 rVec = (float4)(-0.245595,-0.0112916,0.00544129, 0);
24    __constant float stepSize = 0.5;
25    __constant int s = 2;
26    
         if(row < 200 &&  col < 200)  
         {  
                 int i;  
                 float t,x,y,z,probedVal, maxValue = -INFINITY;  
                 float4 t_i, t_j, t_k, value, f, imgPt, pt, v;  
                 int4 n1;  
                 int4 n;  
   
                 float4 d = (float4) (h2[0],h1[0],h1[0],h2[0]);  
                 float4 c = (float4) (h2[1],h1[1],h1[1],h2[1]);  
                 float4 b = (float4) (h2[2],h1[2],h1[2],h2[2]);  
                 float4 a = (float4) (h2[3],h1[3],h1[3],h2[3]);  
   
                 float4 pos = orig + (float)row * rVec + (float)col * cVec;  
                 float4 dir =  (pos - eyeVector) / fabs(pos - eyeVector);  
                 float4 NSize = (float4) ((float)height,(float)length,(float)width,1.0f);  
   
                 pt.w = 1.0f;  
                 pos.w = 1.0f;  
                 dir.w = 1.0f;  
                 value.w = 1.0f;  
27    
28                  for(t = 0.0;  t < 20; t+= stepSize)  __kernel float probe(float * img, int * sAxis, float4 imgPos)
29                  {                  {
30    
31                          pos = pos + stepSize * dir;      float probedVal;
32        float4 f, nf, t, hx, hy, hz;
33                          // Begin Probe Operation      int4 n;
34    
35                          // Transform the value to image space position.      float4 d = (float4) (h[3][0],h[2][0],h[1][0],h[0][0]);
36                          imgPt = (float4) (dot(pos,transformMatrix.s0123),      float4 c = (float4) (h[3][1],h[2][1],h[1][1],h[0][1]);
37                                                            dot(pos,transformMatrix.s4567),      float4 b = (float4) (h[3][2],h[2][2],h[1][2],h[0][2]);
38                                                            dot(pos,transformMatrix.s89ab),      float4 a = (float4) (h[3][3],h[2][3],h[1][3],h[0][3]);
39                                                            dot(pos,transformMatrix.scdef));  
40        f = modf(imgPos,&nf);
41        n = convert_int4(nf);
42    
43    
44        t = (float4) (f.x + 1, f.x, f.x -1, f.x - 2);
45        hx = d + t * (c + t * (b + t*a));
46        t = (float4) (f.y + 1, f.y, f.y - 1, f.y - 2);
47        hy = d + t * (c + t * (b + t*a));
48        t = (float4) (f.z + 1, f.z, f.z - 1, f.z -2);
49        hz = d + t * (c + t * (b + t*a));
50    
51        float vx[4];
52        float vy[4];
53    
54        for(int k = 1-s; k <= s; k++) {             // z is the slowest dimension
55            for(int j = 1-s;  j <= s; j++) {        // y is the medium dimension
56                int index = sAxis[0]*sAxis[1]*(n.z+k) + sAxis[0]*(n.y+j) + (n.x-1);
57                float4 v = (float4)(img[index], img[index+1], img[index+2], img[index+3]);
58                vx[j+s-1] = dot(v,hx);
59            }
60            vy[k+s-1] = dot ((float4) (vx[0],vx[1],vx[2],vx[3]),hy);
61        }
62    
63        probedVal = dot((float4) (vy[0],vy[1],vy[2],vy[3]),hz);
64    
65                          f.xyzw = (float4) (modf(imgPt.x,&x),modf(imgPt.y,&y),modf(imgPt.z,&z),1.0f);      return probedVal;
66    
67                          n1 = (int4) ((int)x,(int)y,(int)z,1.0f);  }
68    
69                          if( (imgPt.x > 1 && imgPt.x < NSize.x - 2) &&  __kernel void mip ( float * img,
70                                  (imgPt.y > 1 && imgPt.y < NSize.y - 2) &&                      float * out,
71                                  (imgPt.z > 1 && imgPt.z < NSize.z - 2))                      float16  transformMatrix,
72                        int2 workDim,
73                        int * sAxis)
74                          {                          {
75    
76        int row = get_global_id(0), col = get_global_id(1);
77    
78        if(row <= workDim.x &&  col <= workDim.y) {
79            float t, probedVal, maxValue = -INFINITY;
80            float4 imgPt;
81    
82                                  // the t value for h(fx - i)          float4 pos = origVec + (float)row * rVec + (float)col * cVec;
83                                  t_i = (float4) (-1.0 - f.x, -f.x,f.x - 1.0, f.x - 2.0);          float4 dir = normalize(pos - eyeVec);
   
                                 // the t value for h(fy - j)  
                                 t_j = (float4) (-1.0 - f.y, -f.y, f.y - 1.0,f.y - 2.0);  
   
                                 // the t value for h(fx - k)  
                                 t_k = (float4) (-1.0 - f.z, -f.z, f.z - 1.0,f.z - 2.0);  
   
   
                                 value = ( (d + t_i * (c + t_i * (b + t_i * a)))  *  // h(fx - i) *  
                                                   (d + t_j * (c + t_j * (b + t_j * a)))  *  // h(fy - j) *  
                                                   (d + t_k * (c + t_k * (b + t_k * a))));   // h(fz - k)  
   
   
                                 n = (int4) (((n1.x-1) * length * width + n1.y-1 * width + n1.z-1),  
                                             (n1.x * length * width +  n1.y * width + n1.z),  
                                             ((n1.x+1) * length * width +  n1.y+1 * width + n1.z+1),  
                                             ((n1.x+2) * length * width + n1.y+2 * width + n1.z+2));  
   
   
   
                                 v = (float4)(img[n.x],  
                                                          img[n.y],  
                                                          img[n.z],  
                                                          img[n.w]);  
84    
85            pos.w = 0.0;
86            dir.w = 0.0;
87    
88                             probedVal  = dot(v,value);  // V(n + <i,j,k>) * summations of the h(x) components          for(t = 0.0;  t < 200; t+= stepSize) {
89                pos = pos + stepSize * dir;
90    
91                // Transform the value to image space position.
92                imgPt = (float4) (dot(pos,transformMatrix.s0123),
93                                  dot(pos,transformMatrix.s4567),
94                                  dot(pos,transformMatrix.s89ab),
95                                  dot(pos,transformMatrix.scdef));
96    
97                            // End Probe Operation              if ((imgPt.x > 1 && imgPt.x < (sAxis[0] - 2)
98                &&  (imgPt.y > 1 && imgPt.y < (sAxis[1] - 2))
99                &&  (imgPt.z > 1 && imgPt.z < (sAxis[2] - 2))) {
100                    probedVal = probe(img,sAxis,imgPt);
101                            if(maxValue < probedVal)                            if(maxValue < probedVal)
102                                          maxValue = probedVal;                                          maxValue = probedVal;
103                          }                          }
104                  }                  }
105                  if(row == 0 && col == 74)          out[row * workDim.x + col] = maxValue;
                         printf("Max Value: %f\n", maxValue );  
                 out[row * 200 + col] = maxValue;  
106          }          }
107  }  }

Legend:
Removed from v.203  
changed lines
  Added in v.441

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