CODEWITHSUNDEEP
opencl
Eduardo
Eduardo

Reputation: 39

OpenCL Optimization

Im new in OpenCL. I wrote an OpenCL kernel to compute grayscale. How Can I optimize that code, is possible? Why the computational time is floating so much? Sometimes Im speedup others not. Im doing something wrong?

kernel code:
kernel void grayscale(__global unsigned char *input)
{
    size_t i = get_global_id(0);    

    float grayscaleValue = (input[i*3] * 0.299F) + (input[i*3+1] * 0.587F) + (input[i*3+2] * 0.114F);

    input[i*3] = grayscaleValue;
    input[i*3+1] = grayscaleValue;
    input[i*3+2] = grayscaleValue;   
}


cpu code:
void GrayScaleCPU(struct PPMFile *ppmStruct)
{    
    for (int i = 0; i < ppmStruct->imageSize; i+=3)
    {
        float greyscaleValue = (ppmStruct->data[i] * 0.299F) + (ppmStruct->data[i+1] * 0.587F) + (ppmStruct->data[i+2] * 0.114F);
        ppmStruct->out[i] = greyscaleValue;
        ppmStruct->out[i+1] = greyscaleValue;
        ppmStruct->out[i+2] = greyscaleValue;
    }
}

int main(void)
{
    struct timespec tS1, tS2;

    tS1.tv_sec = 0;
    tS1.tv_nsec = 0;

    tS2.tv_sec = 0;
    tS2.tv_nsec = 0;

    ...

    clock_settime(CLOCK_REALTIME, &tS1);
    GrayScaleCPU(ppmf);
    clock_gettime(CLOCK_REALTIME, &tS1);

    printf ("Timming took %.12lu seconds to run.\n", tS1.tv_nsec);

    ...

    clock_settime(CLOCK_REALTIME, &tS2);
    GrayScaleOpenCL(ppmf2);
    clock_gettime(CLOCK_REALTIME, &tS2);

    printf ("Timming took %.12lu seconds to run.\n", tS2.tv_nsec);

    float time2 = tS2.tv_nsec;
    float time1 = tS1.tv_nsec;
    float speedup = time2/time1;

    printf ("Speed UP OpenCL/CPU %.20f.\n", speedup);
    return 0;
}

Upvotes: 0

Views: 595

Answers (1)

Austin
Austin

Reputation: 1020

Try buffering your global memory into thread memory:

unsigned char l_input0 = input[i*3];
unsigned char l_input1 = input[i*3 + 1];
unsigned char l_input2 = input[i*3 + 2];
//compute grayscale using l_input0,1,2
input[i*3] = grayscale;
input[i*3 + 1] = grayscale;
input[i*3 + 2] = grayscale;

Also, if your data isn't spaced properly when you call your kernel, you may end up executing on each unsigned char, instead of every 3rd unsigned char as in your for loop example.

You can then go further using local memory and work groups and do your calculations in chunks, though that is more challenging as local work sizes are very device specific and need to be a multiple of the global work size. I've found local work sizes of 16, 32, and 64 work on most devices.

Finally, you benchmarking OpenCL, make sure you are measuring kernel performance and not kernel enqueue time. The easiest way to do this is to start a timer, enqueue you kernel, call clainish on the queue, then stop the timer. There are timing and profiling built into most OpenCL devices which are handled by the queue.

Upvotes: 1

Related Questions