2013-04-10 77 views
4

目前,我有如下就像穿越在OpenCL内核。如果有人对这个相当大的内核进行优化,我会很高兴。的OpenCL内核遍历 - 进一步优化

问题是,我正在SAH BVH上运行此代码,并且我希望在他的论文(了解GPU上的光线遍历效率)中使用遍历的方式获得与Timo Aila类似的性能,当然他的代码使用SplitBVH(我可能会考虑使用它代替SAH BVH,但在我看来它的构建时间确实很慢)。但我问的是遍历,而不是BVH(我迄今为止只适用于场景,SplitBVH不会给你带来比SAH BVH更多的优势)。

首先,这里是我迄今为止(标准的同时,同时遍历内核)。

__constant sampler_t sampler = CLK_FILTER_NEAREST; 

// Inline definition of horizontal max 
inline float max4(float a, float b, float c, float d) 
{ 
    return max(max(max(a, b), c), d); 
} 

// Inline definition of horizontal min 
inline float min4(float a, float b, float c, float d) 
{ 
    return min(min(min(a, b), c), d); 
} 

// Traversal kernel 
__kernel void traverse(__read_only image2d_t nodes, 
         __global const float4* triangles, 
         __global const float4* rays, 
         __global float4* result, 
         const int num, 
         const int w, 
         const int h) 
{ 
    // Ray index 
    int idx = get_global_id(0); 

    if(idx < num) 
    { 
     // Stack 
     int todo[32]; 
     int todoOffset = 0; 

     // Current node 
     int nodeNum = 0; 

     float tmin = 0.0f; 
     float depth = 2e30f; 

     // Fetch ray origin, direction and compute invdirection 
     float4 origin = rays[2 * idx + 0]; 
     float4 direction = rays[2 * idx + 1]; 
     float4 invdir = native_recip(direction); 

     float4 temp = (float4)(0.0f, 0.0f, 0.0f, 1.0f); 

     // Traversal loop 
     while(true) 
     { 
      // Fetch node information 
      int2 nodeCoord = (int2)((nodeNum << 2) % w, (nodeNum << 2)/w); 
      int4 specs = read_imagei(nodes, sampler, nodeCoord + (int2)(3, 0)); 

      // While node isn't leaf 
      while(specs.z == 0) 
      { 
       // Fetch child bounding boxes 
       float4 n0xy = read_imagef(nodes, sampler, nodeCoord); 
       float4 n1xy = read_imagef(nodes, sampler, nodeCoord + (int2)(1, 0)); 
       float4 nz = read_imagef(nodes, sampler, nodeCoord + (int2)(2, 0)); 

       // Test ray against child bounding boxes 
       float oodx = origin.x * invdir.x; 
       float oody = origin.y * invdir.y; 
       float oodz = origin.z * invdir.z; 
       float c0lox = n0xy.x * invdir.x - oodx; 
       float c0hix = n0xy.y * invdir.x - oodx; 
       float c0loy = n0xy.z * invdir.y - oody; 
       float c0hiy = n0xy.w * invdir.y - oody; 
       float c0loz = nz.x * invdir.z - oodz; 
       float c0hiz = nz.y * invdir.z - oodz; 
       float c1loz = nz.z * invdir.z - oodz; 
       float c1hiz = nz.w * invdir.z - oodz; 
       float c0min = max4(min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz), tmin); 
       float c0max = min4(max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz), depth); 
       float c1lox = n1xy.x * invdir.x - oodx; 
       float c1hix = n1xy.y * invdir.x - oodx; 
       float c1loy = n1xy.z * invdir.y - oody; 
       float c1hiy = n1xy.w * invdir.y - oody; 
       float c1min = max4(min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz), tmin); 
       float c1max = min4(max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz), depth); 

       bool traverseChild0 = (c0max >= c0min); 
       bool traverseChild1 = (c1max >= c1min); 

       nodeNum = specs.x; 
       int nodeAbove = specs.y; 

       // We hit just one out of 2 childs 
       if(traverseChild0 != traverseChild1) 
       { 
        if(traverseChild1) 
        { 
         nodeNum = nodeAbove; 
        } 
       } 
       // We hit either both or none 
       else 
       { 
        // If we hit none, pop node from stack (or exit traversal, if stack is empty) 
        if (!traverseChild0) 
        { 
         if(todoOffset == 0) 
         { 
          break; 
         } 
         nodeNum = todo[--todoOffset]; 
        } 
        // If we hit both 
        else 
        { 
         // Sort them (so nearest goes 1st, further 2nd) 
         if(c1min < c0min) 
         { 
          unsigned int tmp = nodeNum; 
          nodeNum = nodeAbove; 
          nodeAbove = tmp; 
         } 

         // Push further on stack 
         todo[todoOffset++] = nodeAbove; 
        } 
       } 

       // Fetch next node information 
       nodeCoord = (int2)((nodeNum << 2) % w, (nodeNum << 2)/w); 
       specs = read_imagei(nodes, sampler, nodeCoord + (int2)(3, 0)); 
      } 

      // If node is leaf & has some primitives 
      if(specs.z > 0) 
      { 
       // Loop through primitives & perform intersection with them (Woop triangles) 
       for(int i = specs.x; i < specs.y; i++) 
       { 
        // Fetch first point from global memory 
        float4 v0 = triangles[i * 4 + 0]; 

        float o_z = v0.w - origin.x * v0.x - origin.y * v0.y - origin.z * v0.z; 
        float i_z = 1.0f/(direction.x * v0.x + direction.y * v0.y + direction.z * v0.z); 
        float t = o_z * i_z; 

        if(t > 0.0f && t < depth) 
        { 
         // Fetch second point from global memory 
         float4 v1 = triangles[i * 4 + 1]; 

         float o_x = v1.w + origin.x * v1.x + origin.y * v1.y + origin.z * v1.z; 
         float d_x = direction.x * v1.x + direction.y * v1.y + direction.z * v1.z; 
         float u = o_x + t * d_x; 

         if(u >= 0.0f && u <= 1.0f) 
         { 
          // Fetch third point from global memory 
          float4 v2 = triangles[i * 4 + 2]; 

          float o_y = v2.w + origin.x * v2.x + origin.y * v2.y + origin.z * v2.z; 
          float d_y = direction.x * v2.x + direction.y * v2.y + direction.z * v2.z; 
          float v = o_y + t * d_y; 

          if(v >= 0.0f && u + v <= 1.0f) 
          { 
           // We got successful hit, store the information 
           depth = t; 
           temp.x = u; 
           temp.y = v; 
           temp.z = t; 
           temp.w = as_float(i); 
          } 
         } 
        } 
       } 
      } 

      // Pop node from stack (if empty, finish traversal) 
      if(todoOffset == 0) 
      { 
       break; 
      } 

      nodeNum = todo[--todoOffset]; 
     } 

     // Store the ray traversal result in global memory 
     result[idx] = temp; 
    } 
} 

当天的第一个问题是,怎么能写在他的OpenCL而持久,而和投机性,同时,而内核?

广告持续而-而,我得到它的权利,我其实刚开始与全球等同于当地的工作尺寸工作尺寸的内核,而这两个数字应该等于以翘曲的GPU /波阵面尺寸? 我得到与CUDA持久线程实现看起来是这样的:

do 
    { 
     volatile int& jobIndexBase = nextJobArray[threadIndex.y]; 

     if(threadIndex.x == 0) 
     { 
       jobIndexBase = atomicAdd(&warpCounter, WARP_SIZE); 
     } 

     index = jobIndexBase + threadIndex.x; 

     if(index >= totalJobs) 
       return; 

     /* Perform work for task numbered 'index' */ 
    } 
    while(true); 

怎么可能相当于OpenCL的样子,我知道我必须做在那里的一些障碍,我也知道,一个后应我原子地将WARP_SIZE添加到warpCounter的分数。

广告投机穿越 - 好,我可能没有任何想法如何这应该在OpenCL的实施,所以任何提示的欢迎。我也不知道在哪里放置障碍物(因为将它们放在模拟__中会导致驾驶员碰撞)。

如果你在这里做到了,感谢阅读&任何提示,答案等,欢迎!

回答

1

你可以做一个优化是使用矢量变量和融合乘法相加功能,以加快建立数学。至于内核的其他部分,它很慢,因为它很分散。如果您可以对信号数据进行假设,则可以通过减少代码分支来缩短执行时间。我没有检查float4 swizles(float 4变量后面的.xxyy和.x .y .z .w),所以检查一下。

  float4 n0xy = read_imagef(nodes, sampler, nodeCoord); 
      float4 n1xy = read_imagef(nodes, sampler, nodeCoord + (int2)(1, 0)); 
      float4 nz = read_imagef(nodes, sampler, nodeCoord + (int2)(2, 0)); 

      float4 oodf4 = -origin * invdir; 

      float4 c0xyf4 = fma(n0xy,invdir.xxyy,oodf4); 

      float4 c0zc1z = fma(nz,(float4)(invdir.z),oodf4); 

      float c0min = max4(min(c0xyf4.x, c0xyf4.y), min(c0xyf4.z, c0xyf4.w), min(c0zc1z.z, c0zc1z.w), tmin); 
      float c0max = min4(max(c0xyf4.x, c0xyf4.y), max(c0xyf4.z, c0xyf4.w), max(c0zc1z.z, c0zc1z.w), depth); 

      float4 c1xy = fma(n1xy,invdir.xxyy,oodf4); 

      float c1min = max4(min(c1xy.x, c1xy.y), min(c1xy.z, c1xy.w), min(c0zc1z.z, c0zc1z.w), tmin); 
      float c1max = min4(max(c1xy.x, c1xy.y), max(c1xy.z, c1xy.w), max(c0zc1z.z, c0zc1z.w), depth);