Pastebin launched a little side project called VERYVIRAL.com, check it out ;-) Want more features on Pastebin? Sign Up, it's FREE!
Guest

An Efficient Parametric Algorithm for Octree Traversal

By: a guest on Oct 30th, 2011  |  syntax: C  |  size: 6.17 KB  |  views: 147  |  expires: Never
download  |  raw  |  embed  |  report abuse  |  print
Text below is selected. Please press Ctrl+C to copy to your clipboard. (⌘+C on Mac)
  1. __device__ void ProcSubtree(Ray & ray,
  2.                                                         float tx0, float ty0, float tz0,
  3.                                                     float tx1, float ty1, float tz1,
  4.                                                     int a, float4 & dstColor)
  5. {
  6.        
  7.         ActualState arrayStack[25];
  8.         int stackTop = 0;
  9.        
  10.         int currNode = 0;
  11.         float txm, tym, tzm;
  12.        
  13.         ActualState state;
  14.         ActualState newState;
  15.  
  16.         state.tx0 = tx0; state.ty0 = ty0; state.tz0 = tz0;
  17.         state.tx1 = tx1; state.ty1 = ty1; state.tz1 = tz1;     
  18.         state.index = 0;
  19.         state.depth = 0;
  20.         arrayStack[stackTop] = state;
  21.        
  22.         float invRGBA = 1.0f / 255.0f;
  23.         uint8 stateLookUp = 0; 
  24.                
  25.        
  26.         while(stackTop >= 0)
  27.         {              
  28.                
  29.                
  30.                 state = arrayStack[stackTop];
  31.                 stackTop--;
  32.                
  33.  
  34.                 if (state.tx1 < 0.0f || state.ty1 < 0.0f || state.tz1 < 0.0f)
  35.                 {
  36.                         continue;
  37.                 }
  38.                        
  39.                 if (state.index < 0)
  40.                 {                      
  41.                         if (state.index != EMPTY_NODE)
  42.                         {      
  43.                                                                
  44.  
  45.                                 //dstColor = make_float4(0);
  46.                                 dstColor.x = 0.0f;
  47.                                 dstColor.y = 0.5f;
  48.                                 dstColor.z = 0.0f;
  49.                                 break;
  50.                         }              
  51.                        
  52.                         continue;
  53.                 }
  54.                
  55.                                
  56.                
  57.                 txm = 0.5f * (state.tx0 + state.tx1);
  58.                 tym = 0.5f * (state.ty0 + state.ty1);
  59.                 tzm = 0.5f * (state.tz0 + state.tz1);
  60.  
  61.                 currNode = FirstNode(state.tx0, state.ty0, state.tz0, txm, tym, tzm);
  62.  
  63.                 state.index++;
  64.                 do
  65.                 {                                              
  66.                         switch (currNode)
  67.                         {
  68.                                 case FLD_NODE:
  69.                                         newState.tx0 = state.tx0; newState.ty0 = state.ty0; newState.tz0 = state.tz0;
  70.                                         newState.tx1 = txm;               newState.ty1 = tym;           newState.tz1 = tzm;                                            
  71.                                        
  72.                                         break;
  73.                                 case BLD_NODE:
  74.                                         newState.tx0 = state.tx0; newState.ty0 = state.ty0; newState.tz0 = tzm;
  75.                                         newState.tx1 = txm;               newState.ty1 = tym;           newState.tz1 = state.tz1;                                                                              
  76.                                        
  77.                                         break;
  78.                                 case FLT_NODE:
  79.                                         newState.tx0 = state.tx0; newState.ty0 = tym;           newState.tz0 = state.tz0;
  80.                                         newState.tx1 = txm;               newState.ty1 = state.ty1; newState.tz1 = tzm;                                                
  81.                                        
  82.                                         break;
  83.                                 case BLT_NODE:
  84.                                         newState.tx0 = state.tx0; newState.ty0 = tym;           newState.tz0 = tzm;
  85.                                         newState.tx1 = txm;               newState.ty1 = state.ty1; newState.tz1 = state.tz1;                                                  
  86.                                        
  87.                                         break;
  88.                                 case FRD_NODE:
  89.                                         newState.tx0 = txm;               newState.ty0 = state.ty0; newState.tz0 = state.tz0;
  90.                                         newState.tx1 = state.tx1; newState.ty1 = tym;           newState.tz1 = tzm;                                            
  91.                                        
  92.                                         break;
  93.                                 case BRD_NODE:
  94.                                         newState.tx0 = txm;               newState.ty0 = state.ty0; newState.tz0 = tzm;
  95.                                         newState.tx1 = state.tx1; newState.ty1 = tym;           newState.tz1 = state.tz1;                                              
  96.                                        
  97.                                         break;
  98.                                 case FRT_NODE:
  99.                                         newState.tx0 = txm;               newState.ty0 = tym;           newState.tz0 = state.tz0;
  100.                                         newState.tx1 = state.tx1; newState.ty1 = state.ty1; newState.tz1 = tzm;                                        
  101.                                        
  102.                                         break;
  103.                                 case BRT_NODE:
  104.                                         newState.tx0 = txm;               newState.ty0 = tym;           newState.tz0 = tzm;
  105.                                         newState.tx1 = state.tx1; newState.ty1 = state.ty1; newState.tz1 = state.tz1;                                          
  106.                                        
  107.                                         break;
  108.                         }
  109.                        
  110.  
  111.                        
  112.                         switch(currNode)
  113.                         {
  114.                                 case FLD_NODE:
  115.                                         stateLookUp = FLD_NODE ^ a;
  116.                                         break;
  117.                                 case BLD_NODE:
  118.                                         stateLookUp = BLD_NODE ^ a;
  119.                                         break;
  120.                                 case FLT_NODE:
  121.                                         stateLookUp = FLT_NODE ^ a;
  122.                                         break;
  123.                                 case BLT_NODE:
  124.                                         stateLookUp = BLT_NODE ^ a;
  125.                                         break;
  126.                                 case FRD_NODE:
  127.                                         stateLookUp = FRD_NODE ^ a;
  128.                                         break;
  129.                                 case BRD_NODE:
  130.                                         stateLookUp = BRD_NODE ^ a;
  131.                                         break;
  132.                                 case FRT_NODE:
  133.                                         stateLookUp = FRT_NODE ^ a;
  134.                                         break;
  135.                                 case BRT_NODE:
  136.                                         stateLookUp = BRT_NODE ^ a;
  137.                                         break;
  138.                                 default:
  139.                                         break;
  140.                         }
  141.  
  142.  
  143.  
  144.                         newState.index  = tex1Dfetch(dataTex, state.index + stateLookUp);
  145.  
  146.                         if (newState.index == 0)
  147.                         {
  148.                                 continue;
  149.                         }
  150.                        
  151.  
  152.  
  153.                         uint8 case1 = 8;
  154.                         uint8 case2 = 8;
  155.                         uint8 case3 = 8;
  156.  
  157.                         switch(currNode)
  158.                         {
  159.                                 case FLD_NODE:
  160.                                         case1 = FRD_NODE; case2 = FLT_NODE; case3 = BLD_NODE;
  161.                                         break;
  162.                                 case BLD_NODE:
  163.                                         case1 = BRD_NODE; case2 = BLT_NODE;
  164.                                         break;
  165.                                 case FLT_NODE:
  166.                                         case1 = FRT_NODE; case3 = BLT_NODE;
  167.                                         break;
  168.                                 case BLT_NODE:
  169.                                         case1 = BRT_NODE;
  170.                                         break;
  171.                                 case FRD_NODE:
  172.                                         case2 = FRT_NODE; case3 = BRD_NODE;
  173.                                         break;
  174.                                 case BRD_NODE:
  175.                                         case2 = BRT_NODE;
  176.                                         break;
  177.                                 case FRT_NODE:
  178.                                         case3 = BRT_NODE;
  179.                                         break;
  180.                                 default:
  181.                                         break;
  182.                         }
  183.  
  184.                         currNode = NewNode(newState.tx1, case1,
  185.                                                        newState.ty1, case2,
  186.                                                            newState.tz1, case3);
  187.  
  188.                        
  189.                         newState.depth = state.depth + 1;
  190.  
  191.                         stackTop++;
  192.                         arrayStack[stackTop] = newState;
  193.                                                                
  194.  
  195.                 } while (currNode < 8);        
  196.         }
  197.                
  198. }
  199.  
  200. __device__ void TraverseRay(Ray & ray, const float3 & dataSize, float4 & dstColor)
  201. {
  202.         int a = 0;
  203.         if (ray.dir.x < 0.0f)
  204.         {                              
  205.                 ray.origin.x = dataSize.x - ray.origin.x;
  206.                 ray.dir.x *= -1;
  207.                 a |= 4;
  208.         }
  209.         if (ray.dir.y < 0.0f)
  210.         {
  211.                 ray.origin.y = dataSize.y - ray.origin.y;
  212.                 ray.dir.y *= -1;
  213.                 a |= 2;
  214.         }
  215.         if (ray.dir.z < 0.0f)
  216.         {
  217.                 ray.origin.z = dataSize.z - ray.origin.z;
  218.                 ray.dir.z *= -1;
  219.                 a |= 1;
  220.         }
  221.  
  222.        
  223.         float invDir = 1.0f / ray.dir.x;
  224.  
  225.         float tx0 = -ray.origin.x * invDir;
  226.         float tx1 = (dataSize.x - ray.origin.x) * invDir;
  227.  
  228.         invDir = 1.0f / ray.dir.y;
  229.         float ty0 = -ray.origin.y * invDir;
  230.         float ty1 = (dataSize.y - ray.origin.y) * invDir;
  231.  
  232.         invDir = 1.0f / ray.dir.z;
  233.         float tz0 = -ray.origin.z * invDir;
  234.         float tz1 = (dataSize.z - ray.origin.z) * invDir;
  235.        
  236.         /*
  237.         float invDir = __fdividef(1.0f, ray.dir.x);
  238.  
  239.         float tx0 = (0 - ray.origin.x) * invDir;
  240.         float tx1 = ( 64 - ray.origin.x) * invDir;
  241.  
  242.         invDir = __fdividef(1.0f, ray.dir.y);
  243.         float ty0 = (0 - ray.origin.y) * invDir;
  244.         float ty1 = ( 64 - ray.origin.y) * invDir;
  245.  
  246.         invDir = __fdividef(1.0f, ray.dir.z);
  247.         float tz0 = (0 - ray.origin.z) * invDir;
  248.         float tz1 = ( 64 - ray.origin.z) * invDir;
  249.         */
  250.  
  251.         if (ray.dir.x == 0.0f)
  252.         {
  253.                 if (tx0 < 0) tx0 = -INT_MAX;
  254.                 else tx0 = INT_MAX;
  255.                 if (tx1 < 0) tx1 = -INT_MAX;
  256.                 else tx1 = INT_MAX;
  257.         }
  258.         if (ray.dir.y == 0.0f)
  259.         {
  260.                 if (ty0 < 0) ty0 = -INT_MAX;
  261.                 else ty0 = INT_MAX;
  262.                 if (ty1 < 0) ty1 = -INT_MAX;
  263.                 else ty1 = INT_MAX;
  264.         }
  265.         if (ray.dir.z == 0.0f)
  266.         {
  267.                 if (tz0 < 0) tz0 = -INT_MAX;
  268.                 else tz0 = INT_MAX;
  269.                 if (tz1 < 0) tz1 = -INT_MAX;
  270.                 else tz1 = INT_MAX;
  271.         }
  272.  
  273.         if (Max(tx0, ty0, tz0) < Min(tx1, ty1, tz1))
  274.         {                              
  275.                 //dstColor.x = 0.5f;
  276.                 ProcSubtree(ray, tx0, ty0, tz0, tx1, ty1, tz1, a, dstColor);                                                                                           
  277.         }
  278.        
  279.  
  280. }
  281.  
  282.  
  283.