OpenCL implementations of IQZZ and IDCT for MJPEG

Question

I am using this code for MJPEG decoding and I am trying to make two functions (IQZZ and IDCT) run faster on the GPU (NVIDIA Tesla k20c). I am using the OpenCL framework to accomplish this task.

I have already successfully offloaded these functions to the GPU and am getting the expected output. However, the output video is very slow after offloading the code to the GPU.

My .cl file is as follows:

/******************************* IDCT *************************************/

void idct_1D(__local int *Y);

__kernel void IDCT(__global int* input, __global uchar* output) 
{
unsigned int kid= get_global_id(0);

    __local int Y[64]; 
    int k= get_global_id(0);
    int l;
    int lid= get_global_id(1);
    __local int Yc[8];

   if (k < 8)
    {
        for (l = 0; l < 8; l++) 

     {
     Y(k, l) = SCALE(input[(k << 3) + l], S_BITS);
     }
        idct_1D(&Y(k, 0));
    }

    if (lid < 8)
    {

        for (k = 0; k < 8; k++)
    {
            Yc[k] = Y(k, lid);
    }

        idct_1D(Yc);

        for (k = 0; k < 8; k++)
        {
            int r = 128 + DESCALE(Yc[k], S_BITS + 3);
            r = r > 0 ? (r < 255 ? r : 255) : 0;
            X(k, lid) = r;
        }

    }
}

void idct_1D(__local int *Y) 
{

int z1[8], z2[8], z3[8];


    but(Y[0], Y[4], z1[1], z1[0]);
    rot(1, 6, Y[2], Y[6], &z1[2], &z1[3]);
    but(Y[1], Y[7], z1[4], z1[7]);
    z1[5] = CMUL(sqrt2, Y[3]);
    z1[6] = CMUL(sqrt2, Y[5]);

    but(z1[0], z1[3], z2[3], z2[0]);
    but(z1[1], z1[2], z2[2], z2[1]);
    but(z1[4], z1[6], z2[6], z2[4]);
    but(z1[7], z1[5], z2[5], z2[7]);

    z3[0] = z2[0];
    z3[1] = z2[1];
    z3[2] = z2[2];
    z3[3] = z2[3];
    rot(0, 3, z2[4], z2[7], &z3[4], &z3[7]);
    rot(0, 1, z2[5], z2[6], &z3[5], &z3[6]);

    but(z3[0], z3[7], Y[7], Y[0]);
    but(z3[1], z3[6], Y[6], Y[1]);
    but(z3[2], z3[5], Y[5], Y[2]);
    but(z3[3], z3[4], Y[4], Y[3]);
} 

/*---------------IQZZ----------------------------*/

__kernel void iqzz_block(__global int in[64], __global int out[64],
        __global uchar table[64])
{
    uint index= get_global_id(0);
    int priv_in[64];
    uchar priv_table[64];

    int priv_out[64];


    if (index < 64)
    {   

    priv_in[index]= in[index];

    priv_table[index]= table[index];

        priv_out[G_ZZ[index]] = priv_in[index] * priv_table[index];

        out[G_ZZ[index]]= priv_out[G_ZZ[index]];

    }
}

For IDCT, I simply copied and pasted constants from the the .c file. I haven't included the constants in my query for conciseness. Details of the constants can be found here.

In main.c, I have simply substituted the function calls with OpenCL commands to transfer data to the device, execute the kernel there and transmit the results back on the CPU.

My main.c looks like this:

/* Get Platform */
ret= clGetPlatformIDs(1, &platform_id, &ret_num_platforms);

/* Get Device */
ret= clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &ret_num_devices);

 /* Create Context */
context = clCreateContext(0, 1, &device_id, NULL, NULL, &ret);

/* Create Command Queue */
command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

/* Create kernel from source */
program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);


ret= clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

//--------kernel for iqzz-----------//
kernel= clCreateKernel(program, "iqzz_block", &ret);

//-------kernel for idct-----------//
cos_kernel= clCreateKernel(program, "IDCT", &ret);

cl_mem block_GPU = clCreateBuffer(context, CL_MEM_READ_WRITE, 64 * sizeof(cl_int), NULL, &ret);


//This will serve as the output buffer for iqzz
cl_mem DCT_Input = clCreateBuffer(context, CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR, 64 * sizeof(cl_int), unZZ_MCU, &ret);
chk(ret, "clCreateBuffer");

//Output buffer
cl_mem  DCT_Output = clCreateBuffer(context, CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR, (MCU_sx * MCU_sy * max_ss_h * max_ss_v) + 4, YCbCr_MCU_ds[component_index] + (64 * chroma_ss), &ret);

//Regular code from main.c follows............

case M_SOS:

//regular code from main.c.......

//The Relevant part starts here......

for (index_X = 0; index_X < nb_MCU_X; index_X++) {

for (index_Y = 0; index_Y < nb_MCU_Y; index_Y++) {

for (index = 0; index < SOS_section.n; index++)
{

int component_index = component_order[index];

int nb_MCU = ((SOF_component[component_index].HV>> 4) & 0xf)*(SOF_component[component_index].HV & 0x0f);

for (chroma_ss = 0; chroma_ss < nb_MCU; chroma_ss++)
{
unpack_block(movie, &scan_desc,index, MCU);

/////--------------Transfer data to buffers----------------////////////

ret = clEnqueueWriteBuffer(command_queue, block_GPU, CL_TRUE, 0, 64 * sizeof(cl_int), MCU, 0, NULL, NULL);

ret = clEnqueueWriteBuffer(command_queue, qtable_GPU, CL_TRUE, 0, 64 * sizeof(cl_uchar), DQT_table[SOF_component[component_index].q_table], 0, NULL, NULL);

cl_mem qtable_GPU = clCreateBuffer(context, CL_MEM_READ_WRITE, 64 * sizeof(cl_uchar), NULL, &ret);

/* Set OpenCL kernel arguments */
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&block_GPU);
ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&DCT_Input);
ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&qtable_GPU);

start_time = wtime();

size_t global=64;
size_t local= 16;

ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, &local, 0, NULL, NULL);

run_time += wtime() - start_time;

//Copy result from device to host
ret = clEnqueueReadBuffer(command_queue, DCT_Input, CL_TRUE, 0, 64 * sizeof(cl_int), &unZZ_MCU, 0, NULL, NULL);

/////---------------IDCT-----------------//////

ret = clSetKernelArg(cos_kernel, 0, sizeof(cl_mem), (void *)&DCT_Input);
ret |= clSetKernelArg(cos_kernel, 1, sizeof(cl_mem), (void *)&DCT_Output);

//No. of work-items
const size_t globalForInverseDCT[2]= {8, 8};

ret = clEnqueueNDRangeKernel(command_queue, cos_kernel, 1, NULL, &globalForInverseDCT, &localForInverseDCT, 0, NULL, NULL);

ret = clEnqueueReadBuffer(command_queue, DCT_Output, CL_TRUE, 0, (MCU_sx * MCU_sy * max_ss_h * max_ss_v) + 4, YCbCr_MCU_ds[component_index] + (64 * chroma_ss), 0, NULL, NULL);

}                       
upsampler(YCbCr_MCU_ds[component_index],YCbCr_MCU[component_index],Horizontal,Vertical,max_ss_h,max_ss_v);
}

//more code which is not immediately relevant follows......

}

How can I modify my iqzz and idct kernels to make them run faster on the GPU?

The details of my GPU are as follows:

DEVICE_NAME = Tesla K20c
DEVICE_VENDOR = NVIDIA Corporation
DEVICE_VERSION = OpenCL 1.2 CUDA
DRIVER_VERSION = 352.21
DEVICE_MAX_COMPUTE_UNITS = 13
DEVICE_MAX_CLOCK_FREQUENCY = 705
DEVICE_GLOBAL_MEM_SIZE = 5032706048
CL_DEVICE_ERROR_CORRECTION_SUPPORT: yes
CL_DEVICE_LOCAL_MEM_TYPE:       local
CL_DEVICE_LOCAL_MEM_SIZE:       48 KByte
CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE: 64 KByte
CL_DEVICE_QUEUE_PROPERTIES:     CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
CL_DEVICE_QUEUE_PROPERTIES:     CL_QUEUE_PROFILING_ENABLE

Answer 1

1. Indent your loop bodies.
2. Actually check ret - you're uselessly assigning and discarding it every time.
3. Use better variable names: avoid single letters (Y, k, l) and generic names (index)
4. It appears all the work of your code is inside four nested loops:
   1. Try to vectorise - rewrite the inner block to operate on muliple pixels/components/chromas simultaneously.
   2. Optimise the iteration order for caching and branch prediction.
   3. Pull up anything that doesn't actually have to be in the inner loop (like the runtime calculations?)

The main problem with offloading to a GPU is that the data transfer is very slow. You need to minimise the number of copies to/from the GPU and take maximum advantage of its parallelism. If you cannot do that then it will always be faster to stay on the CPU (especially with SIMD); even if the individual operations are slower it can still get through the data quicker.