13
\$\begingroup\$

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
\$\endgroup\$
  • \$\begingroup\$ Why haven't I received any feedback for my question? Do I need to change anything to get responses? \$\endgroup\$ – a_sid Sep 9 '17 at 19:25
  • 4
    \$\begingroup\$ The question is really difficult. I don't think anyone will be able to improve the performance just looking at the code without access to the appropriate environment in which to profile. You're the only person with the ability to find the bottlenecks. \$\endgroup\$ – Stop Harming Monica Nov 1 '17 at 23:35
  • 9
    \$\begingroup\$ @a_sid You are not getting comments because your code is not written in a way that others can read it. You say you have a performance issue but not what or where. You have a method iqzz_block() that you appear to want reviewed but assign it to the name kernel which is referenced 22 times. If you couldn't make the effort to comment or even indent your code, why do you expect anyone to take the effort to decode it? Spend a bit of time writing python or some other 'read_many' language, it will make your c more intelligible. Sorry, but I down voted for those reasons. \$\endgroup\$ – Paul Smith Nov 2 '17 at 12:14
14
+400
\$\begingroup\$
  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.

\$\endgroup\$
  • 3
    \$\begingroup\$ This is very general (and very good) advice on how to compute efficiently on the GPU. Probably the maximum one can say to the question without spending much more time. \$\endgroup\$ – Trilarion Nov 3 '17 at 10:57

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Not the answer you're looking for? Browse other questions tagged or ask your own question.