优化内核混洗键代码 - OpenCL

Question

我刚刚开始接触 OpenCL 并了解编写内核代码的基础知识。我已经编写了一个内核代码来计算点数组的洗牌键。因此，对于许多点 N，混洗后的密钥以 3 位方式计算，其中深度 d 处的 x 位（0

xd = 0 if p.x < Cd.x  
xd = 1, otherwise

Shuffled xyz 键为：

x1y1z1x2y2z2...xDyDzD 

编写的内核代码如下。该点以列主要格式输入。

__constant float3 boundsOffsetTable[8] = {
              {-0.5,-0.5,-0.5},
              {+0.5,-0.5,-0.5},
              {-0.5,+0.5,-0.5},
              {-0.5,-0.5,+0.5},
              {+0.5,+0.5,-0.5},
              {+0.5,-0.5,+0.5},
              {-0.5,+0.5,+0.5},
              {+0.5,+0.5,+0.5}
};
uint setBit(uint x,unsigned char position)
{
uint mask = 1<<position;
return x|mask;
}

__kernel void morton_code(__global float* point,__global uint*code,int level, float3          center,float radius,int size){
// Get the index of the current element to be processed
int i = get_global_id(0);
float3 pt; 
pt.x = point[i];pt.y = point[size+i]; pt.z = point[2*size+i];
code[i] = 0;
float3 newCenter;
float newRadius;
if(pt.x>center.x) code = setBit(code,0);
if(pt.y>center.y) code = setBit(code,1);
if(pt.z>center.z) code = setBit(code,2);
for(int l = 1;l<level;l++)
{
    for(int i=0;i<8;i++)
    {
        newRadius = radius *0.5;
        newCenter = center + boundOffsetTable[i]*radius;
        if(newCenter.x-newRadius<pt.x && newCenter.x+newRadius>pt.x && newCenter.y-newRadius<pt.y && newCenter.y+newRadius>pt.y && newCenter.z-newRadius<pt.z && newCenter.z+newRadius>pt.z)
        {
            if(pt.x>newCenter.x) code = setBit(code,3*l);
            if(pt.y>newCenter.y) code = setBit(code,3*l+1);
            if(pt.z>newCenter.z) code = setBit(code,3*l+2);
        }
    }
}
}

它有效，但我只是想问一下我是否在代码中遗漏了什么，以及是否有办法优化代码。

Answer 1

试试这个内核：

__kernel void morton_code(__global float* point,__global uint*code,int level, float3          center,float radius,int size){
// Get the index of the current element to be processed
int i = get_global_id(0);
float3 pt; 
pt.x = point[i];pt.y = point[size+i]; pt.z = point[2*size+i];
uint res;
res = 0;
float3 newCenter;
float newRadius;
if(pt.x>center.x) res = setBit(res,0);
if(pt.y>center.y) res = setBit(res,1);
if(pt.z>center.z) res = setBit(res,2);
for(int l = 1;l<level;l++)
{
    for(int i=0;i<8;i++)
    {
        newRadius = radius *0.5;
        newCenter = center + boundOffsetTable[i]*radius;
        if(newCenter.x-newRadius<pt.x && newCenter.x+newRadius>pt.x && newCenter.y-newRadius<pt.y && newCenter.y+newRadius>pt.y && newCenter.z-newRadius<pt.z && newCenter.z+newRadius>pt.z)
        {
            if(pt.x>newCenter.x) res = setBit(res,3*l);
            if(pt.y>newCenter.y) res = setBit(res,3*l+1);
            if(pt.z>newCenter.z) res = setBit(res,3*l+2);
        }
    }
}
//Save the result
code[i] = res;
}

要优化的规则：

1. 避免使用全局内存（您直接使用全局内存中的“代码”，我更改了它），您现在应该会看到性能提高了 3 倍。
2. 避免使用 If，如果可能，请改用“select”。 （参见 OpenCL 文档）
3. 在内核中使用更多内存。您不需要在位级别进行操作。 int 级别的操作会更好，并且可以避免对“setBit”的大量调用。然后你可以在最后构建你的结果。

另一个有趣的事情。如果您在 3D 级别操作，您可以使用 float3 变量并使用 OpenCL 运算符计算距离。这可以大大提高你的表现。 BUt 还需要完全重写您的内核。