OpenCL Demo1测试与性能评估

前言

通过一个openCL的测试demo来分析评估高通芯片的GPU运算性能

一、芯片平台

Hardware : Qualcomm Technologies, Inc SDM660

$ cat /proc/cpuinfo
Processor       : AArch64 Processor rev 2 (aarch64)
processor       : 0
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x801
CPU revision    : 4

processor       : 1
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x801
CPU revision    : 4

processor       : 2
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x801
CPU revision    : 4

processor       : 3
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x801
CPU revision    : 4

processor       : 4
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x800
CPU revision    : 2

processor       : 5
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x800
CPU revision    : 2

processor       : 6
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x800
CPU revision    : 2

processor       : 7
BogoMIPS        : 38.40
Features        : fp asimd evtstrm aes pmull sha1 sha2 crc32
CPU implementer : 0x51
CPU architecture: 8
CPU variant     : 0xa
CPU part        : 0x800
CPU revision    : 2

Hardware        : Qualcomm Technologies, Inc SDM660

GPU型号：adreno512

二、测试代码

int kerneltest()
{
    cl_device_id device;
    cl_context context;
    cl_command_queue command_queue;
    cl_program program;
    cl_kernel kernel;
    cl_mem buffer_src;
    cl_mem buffer_dst;
    cl_int err_num = CL_SUCCESS;
    cl_uint buffer_size_in_bytes;
    timeval start;

    context = CreateContext(&device);
    if (NULL == context)
    {
        printf("MainError:Create Context Failed!\n");
        return -1;
    }

    command_queue = CreateCommandQueue(context, device);
    if (NULL == command_queue)
    {
        printf("MainError:Create CommandQueue Failed!\n");
        return -1;
    }

    char *device_source_str = ClUtilReadFileToString("/storage/emulated/0/test/kerneltest.cl");
    program                 = CreateProgram(context, device, device_source_str);
    if (NULL == program)
    {
        printf("MainError:Create CommandQueue Failed!\n");
        return -1;
    }

    kernel = CreateKernel(program, "TransposeKernel", device);
    if (NULL == kernel)
    {
        printf("MainError:Create Kernel Failed!\n");
        return -1;
    }

    const int c_loop_count = 30;

    int width                          = 4096;
    int height                         = 4096;//4096;
    buffer_size_in_bytes               = width * height * sizeof(cl_uchar);
    cl_uchar *host_src_matrix          = (cl_uchar *)malloc(buffer_size_in_bytes);
    cl_uchar *host_transposed_matrix   = (cl_uchar *)malloc(buffer_size_in_bytes);
    cl_uchar *device_transposed_matrix = (cl_uchar *)malloc(buffer_size_in_bytes);
    memset(device_transposed_matrix, 0, buffer_size_in_bytes);
    DataInit(host_src_matrix, width, height);
    printf("Matrix Width =%d Height=%d\n", width, height);
    gettimeofday(&start, NULL);
    for (int i = 0; i < c_loop_count; i++)
    {
        CpuTranspose(host_src_matrix, host_transposed_matrix, width, height);
    }
    PrintDuration(&start, "Cpu Transpose", c_loop_count);

    buffer_src = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                buffer_size_in_bytes, host_src_matrix, &err_num);
    CheckClStatus(err_num, "Create src buffer");
    buffer_dst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, buffer_size_in_bytes, NULL, &err_num);
    CheckClStatus(err_num, "Create dst buffer");

    err_num = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_src);
    err_num |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer_dst);
    err_num |= clSetKernelArg(kernel, 2, sizeof(int), &width);
    err_num |= clSetKernelArg(kernel, 3, sizeof(int), &height);
    CheckClStatus(err_num, "Set Kernel Arg");

    size_t global_work_size[3];
    size_t local_work_size[3];
    local_work_size[0] = 32;
    local_work_size[1] = 32;//32*32 = 1024
    local_work_size[2] = 0;

    global_work_size[0] =
        (width + local_work_size[0] - 1) / local_work_size[0] * local_work_size[0];
    global_work_size[1] =
        (height + local_work_size[1] - 1) / local_work_size[1] * local_work_size[1];
    global_work_size[2] = 0;

    printf("global_work_size=(%zu,%zu)\n", global_work_size[0], global_work_size[1]);

    cl_event kernel_event = NULL;
    //------------------------------------------

        printf("--------------------------------------------------------\n");
        gettimeofday(&start, NULL);
        for (int i = 0; i < c_loop_count; i++)
        {
            err_num = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work_size,
                                             NULL, 0, NULL, &kernel_event);
            CheckClStatus(err_num, "ClEnqueueNDRangeKernel");
            err_num = clWaitForEvents(1, &kernel_event);
            CheckClStatus(err_num, "ClWaitForEvents");
        }
        PrintDuration(&start, "OpenCL Transpose", c_loop_count);
        
        printf("--------------------------------------------------------\n");
        printf("local_work_size=(%zu,%zu)\n", local_work_size[0], local_work_size[1]);
        gettimeofday(&start, NULL);
        for (int i = 0; i < c_loop_count; i++)
        {
            err_num = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work_size,
                                             local_work_size, 0, NULL, &kernel_event);
            CheckClStatus(err_num, "ClEnqueueNDRangeKernel");
            err_num = clWaitForEvents(1, &kernel_event);
            CheckClStatus(err_num, "ClWaitForEvents");
        }
        PrintDuration(&start, "OpenCL Transpose", c_loop_count);
        
    //PrintProfilingInfo(kernel_event);

    err_num = clEnqueueReadBuffer(command_queue, buffer_dst, CL_TRUE, 0, buffer_size_in_bytes,
                                  device_transposed_matrix, 0, NULL, NULL);

    DataCompare(host_transposed_matrix, device_transposed_matrix, width, height);

    free(device_source_str);
    free(host_src_matrix);
    free(host_transposed_matrix);
    free(device_transposed_matrix);
    clReleaseEvent(kernel_event);
    clReleaseMemObject(buffer_src);
    clReleaseMemObject(buffer_dst);

    CleanUp(context, command_queue, program, kernel);
    return 0;
}

三、测试分析

3.1 demo运算性能评估

 D/PRINTJ: max_compute_units = 2
 D/PRINTJ: global_mem_size = 1959766016
 D/PRINTJ: max_work_item_size_of_work_group_dim 0=1024
 D/PRINTJ: max_work_item_size_of_work_group_dim 1=1024
 D/PRINTJ: max_work_item_size_of_work_group_dim 2=1024
 D/PRINTJ: Kernel TransposeKernel max workgroup size=1024
 D/PRINTJ: Kernel TransposeKernel perferred workgroup size multiple=1024
 D/PRINTJ: Matrix Width =4096 Height=4096
 D/PRINTJ: Cpu Transpose consume average time: 523958 us
 D/PRINTJ: global_work_size=(4096,4096)
 D/PRINTJ: --------------------------------------------------------
 D/PRINTJ: OpenCL Transpose consume average time: 211879 us
 D/PRINTJ: --------------------------------------------------------
 D/PRINTJ: local_work_size=(32,32)
 D/PRINTJ: OpenCL Transpose consume average time: 34605 us
 D/PRINTJ: A and B match!

通过上述测试发现设置 local_work_size = {32，32};  2*32 = 1024，GPU运算性能可以获得极大提升，加速比达到523958/34605 = 15.14。