Optimizing kernel shuffled keys code - OpenCL

Question

I have just started getting into OpenCL and going through the basics of writing a kernel code. I have written a kernel code for calculating shuffled keys for points array. So, for a number of points N, the shuffled keys are calculated in 3-bit fashion, where x-bit at depth d (0

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

The Shuffled xyz key is given as:

x1y1z1x2y2z2...xDyDzD

The Kernel code written is given below. The point is inputted in a column major format.

__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);
        }
    }
}
}

It works but I just wanted to ask if I am missing something in the code and if there is an way to optimize the code.

What's the current execution time, how much faster would you like your code to be, and do you want to make your code platform specific or not? — ardiyu07

DarkZeros DarkZeros · Accepted Answer · 2013-05-15T12:17:01

Try this kernel:

__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;
}

Rules to optimize:

Avoid Global memory (you were using "code" directly from global memory, I changed that), you should see 3x increase in performance now.
Avoid Ifs, use "select" instead if it is possible. (See OpenCL documentation)
Use more memory inside the kernel. You don't need to operate at bit level. Operation at int level would be better and could avoid huge amount of calls to "setBit". Then you can construct your result at the end.

Another interesting thing. Is that if you are operating at 3D level, you can just use float3 variables and compute the distances with OpenCL operators. This can increase your performance quite a LOT. BUt also requires a complete rewrite of your kernel.

Optimizing kernel shuffled keys code - OpenCL

1 Answers