cuda - 内核优化答案

【问题标题】：cuda - kernel optimizationcuda - 内核优化
【发布时间】：2011-10-22 09:43:47
【问题描述】：

我创建了一个简单的粒子系统。我有一台具有计算能力 2.1 的设备。我可以改变什么来优化内核？

我假设变量tPos 和tVel 存储在寄存器中。

__global__ void particles_kernel(float4 *vbo, float4 *pos, float4 *vel)
{
     int tid = blockIdx.x * blockDim.x + threadIdx.x;

     float4 tPos = pos[tid];
     float4 tVel = vel[tid];

     tPos.x += tVel.x;
     tPos.y += tVel.y;
     tPos.z += tVel.z;

     if(tPos.x < -2.0f)
     {
         tVel.x = -tVel.x;
     }
     else if(tPos.x > 2.0f)
     {
         tVel.x = -tVel.x;
     }


     if(tPos.y < -2.0f)
     {
         tVel.y = -tVel.y;
     }
     else if(tPos.y > 2.0f)
     {
         tVel.y = -tVel.y;
     }


     if(tPos.z < -2.0f)
     {
         tVel.z = -tVel.z;
     }
     else if(tPos.z > 2.0f)
     {
         tVel.z = -tVel.z;
     }


     pos[tid] = tPos;
     vel[tid] = tVel;


     vbo[tid] = make_float4(tPos.x, tPos.y, tPos.z, tPos.w);
}

【问题讨论】：

如果不笼统地回答，这将是非常困难的。您可能想查看“Optimizing CUDA part I (PDF)”和“Optimizing CUDA part II (PDF)”以获取一些一般提示和技巧。其他优秀资源是“Programming Massively Parallel Processors”和“CUDA by Example”。

标签： cuda nvidia gpu

【解决方案1】：

除非我遗漏了什么，你的钳位代码可以这样简化：

if (fabsf(tVel.x) > 2.0f) tVel.x = -tVel.x;
if (fabsf(tVel.y) > 2.0f) tVel.y = -tVel.y;
if (fabsf(tVel.z) > 2.0f) tVel.z = -tVel.z;

但是，鉴于计算量相对较小，此更改可能不会提高性能，因为代码似乎受内存限制（您正在通过数据流式传输）。也许您的应用程序中的其他地方有额外的计算，您可以将这些计算与此计算结合起来以增加计算密度？

【讨论】：

是的，这个内核非常简单。您可能希望在一个循环中为每个线程处理多个粒子——在计算这几个粒子之前先加载几个粒子，以便您有更多的指令与内存访问和更多的内存事务重叠。否则，这里没有太多可优化的地方。