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 302, Tue Aug 17 03:52:21 2010 UTC revision 441, Wed Oct 20 19:21:32 2010 UTC
# Line 1  Line 1 
1  /* mip.cl  /* mip.cl
2   *   *
3   * COPYRIGHT (c) 2010 The Diderot Project (http://diderot.cs.uchicago.edu)   * COPYRIGHT (c) 2010 The Diderot Project (http://diderot-language.cs.uchicago.edu)
4   * All rights reserved.   * All rights reserved.
5   *   *
6   * An OpenCL file containing a probe kernel and mip kernel implementation   * An OpenCL file containing a probe kernel and mip kernel implementation
# Line 13  Line 13 
13          { 1.33333,  -2.0,  1.0, -0.166667 },   //  1 .. 2          { 1.33333,  -2.0,  1.0, -0.166667 },   //  1 .. 2
14  };  };
15    
16  __constant float4 eyeVec = (float4)(25,15,10,1);  //__constant float4 eyeVec = (float4)(25,15,10,1);
17  __constant float4 origVec = (float4)(8.83877,2.5911,7.65275,0);  //__constant float4 origVec = (float4)(8.83877,2.5911,7.65275,0);
18  __constant float4 cVec= (float4)(-0.0151831,0.0278357,0,0);  //__constant float4 cVec= (float4)(-0.0151831,0.0278357,0,0);
19  __constant float4 rVec = (float4)(0.0074887,0.00408474,-0.0305383,0);  //__constant float4 rVec = (float4)(0.0074887,0.00408474,-0.0305383,0);
20  __constant float stepSize = 0.1;  __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;  __constant int s = 2;
26    
27    
28  __kernel float probe(float * img, int * sAxis, float4 imgPos) {  __kernel float probe(float * img, int * sAxis, float4 imgPos)
29    {
30    
31         float probedVal;         float probedVal;
32         float4 f, nf, t, hx, hy, hz;         float4 f, nf, t, hx, hy, hz;
# Line 43  Line 48 
48          t = (float4) (f.z + 1, f.z, f.z - 1, f.z -2);          t = (float4) (f.z + 1, f.z, f.z - 1, f.z -2);
49          hz = d + t * (c + t * (b + t*a));          hz = d + t * (c + t * (b + t*a));
50    
51          float vy[4] = {0.0f,0.0f,0.0f,0.0f};      float vx[4];
52          float vz[4] = {0.0f,0.0f,0.0f,0.0f};      float vy[4];
53    
54          for(int i = 1-s; i <= s; i++) {      for(int k = 1-s; k <= s; k++) {             // z is the slowest dimension
55           for(int j = 1-s;  j <= s; j++) {          for(int j = 1-s;  j <= s; j++) {        // y is the medium dimension
56              float4 v = (float4) (              int index = sAxis[0]*sAxis[1]*(n.z+k) + sAxis[0]*(n.y+j) + (n.x-1);
57                       img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z-1))],              float4 v = (float4)(img[index], img[index+1], img[index+2], img[index+3]);
58                       img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z))],              vx[j+s-1] = dot(v,hx);
                      img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z+1))],  
                      img[((int)n.x+i) + sAxis[0]*(((int)n.y+j) + sAxis[1]*((int)n.z+2))]);  
             vz[j+s-1] += dot(v,hz);  
59           }           }
60                  vy[i+s-1] = dot ((float4) (vz[0],vz[1],vz[2],vz[3]),hy);          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]),hx);      probedVal = dot((float4) (vy[0],vy[1],vy[2],vy[3]),hz);
64    
65          return probedVal;          return probedVal;
66    
67  }  }
68    
69  __kernel void mip ( float * img,  __kernel void mip ( float * img,
70                                          float * out,                                          float * out,
71                                          float16  transformMatrix,                                          float16  transformMatrix,
72                                          int2 workDim,                                          int2 workDim,
73                                          int * sAxis){                      int * sAxis)
74    {
75    
76          int row = get_global_id(0), col = get_global_id(1);          int row = get_global_id(0), col = get_global_id(1);
77    
78          if(row <= workDim.x &&  col <= workDim.y)      if(row <= workDim.x &&  col <= workDim.y) {
         {  
79                  float t, probedVal, maxValue = -INFINITY;                  float t, probedVal, maxValue = -INFINITY;
80                  float4 imgPt;                  float4 imgPt;
81    
# Line 82  Line 85 
85                  pos.w = 0.0;                  pos.w = 0.0;
86                  dir.w = 0.0;                  dir.w = 0.0;
87    
88                  for(t = 0.0;  t < 40; t+= stepSize)          for(t = 0.0;  t < 200; t+= stepSize) {
                 {  
89                          pos = pos + stepSize * dir;                          pos = pos + stepSize * dir;
90    
91                          // Transform the value to image space position.                          // Transform the value to image space position.
# Line 92  Line 94 
94                                            dot(pos,transformMatrix.s89ab),                                            dot(pos,transformMatrix.s89ab),
95                                            dot(pos,transformMatrix.scdef));                                            dot(pos,transformMatrix.scdef));
96    
97                if ((imgPt.x > 1 && imgPt.x < (sAxis[0] - 2)
98                          if( (imgPt.x > 1 && imgPt.x < (sAxis[0] - 2)) &&              &&  (imgPt.y > 1 && imgPt.y < (sAxis[1] - 2))
99                                  (imgPt.y > 1 && imgPt.y < (sAxis[1]- 2)) &&              &&  (imgPt.z > 1 && imgPt.z < (sAxis[2] - 2))) {
                                 (imgPt.z > 1 && imgPt.z < (sAxis[2] - 2)))  
                         {  
100                                  probedVal = probe(img,sAxis,imgPt);                                  probedVal = probe(img,sAxis,imgPt);
101                                  if(maxValue < probedVal)                                  if(maxValue < probedVal)
102                                          maxValue = probedVal;                                          maxValue = probedVal;

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

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