logo标签列表

C vs OpenCL，如何比较时间测量结果？

cperformancealgorithmopenclmeasurement
3

在之前的帖子中，我提出了关于C时间测量的问题。现在，我想知道如何比较C“函数”和OpenCL“函数”的结果。

这是主机OpenCL和C的代码：

#define PROGRAM_FILE "sum.cl"
#define KERNEL_FUNC "float_sum"
#define ARRAY_SIZE 1000000


#include <stdio.h>
#include <stdlib.h>
#include <time.h>

#include <CL/cl.h>

int main()
{
    /* OpenCL Data structures */

    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_program program;
    cl_kernel kernel;    
    cl_command_queue queue;
    cl_mem vec_buffer, result_buffer;

    cl_event prof_event;;

    /* ********************* */

    /* C Data Structures / Data types */
    FILE *program_handle; //Kernel file handle
    char *program_buffer; //Kernel buffer

    float *vec, *non_parallel;
    float result[ARRAY_SIZE];

    size_t program_size; //Kernel file size

    cl_ulong time_start, time_end, total_time;

    int i;
    /* ****************************** */

    /* Errors */
    cl_int err;
    /* ****** */

    non_parallel = (float*)malloc(ARRAY_SIZE * sizeof(float));
    vec          = (float*)malloc(ARRAY_SIZE * sizeof(float));

    //Initialize the vector of floats
    for(i = 0; i < ARRAY_SIZE; i++)
    vec[i] = i + 1;

    /************************* C Function **************************************/
    clock_t start, end;

    start = clock();

    for( i = 0; i < ARRAY_SIZE; i++) 
    {
    non_parallel[i] = vec[i] * vec[i];
    }
    end = clock();
    printf( "Number of seconds: %f\n", (clock()-start)/(double)CLOCKS_PER_SEC );

    free(non_parallel);
    /***************************************************************************/




    clGetPlatformIDs(1, &platform, NULL);//Just want NVIDIA platform
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);

    // Context error?
    if(err)
    {
    perror("Cannot create context");
    return 1;
    }

    //Read the kernel file
    program_handle = fopen(PROGRAM_FILE,"r");
    fseek(program_handle, 0, SEEK_END);
    program_size = ftell(program_handle);
    rewind(program_handle);

    program_buffer = (char*)malloc(program_size + 1);
    program_buffer[program_size] = '\0';
    fread(program_buffer, sizeof(char), program_size, program_handle);
    fclose(program_handle);

    //Create the program
    program = clCreateProgramWithSource(context, 1, (const char**)&program_buffer, 
                    &program_size, &err);

    if(err)
    {
    perror("Cannot create program");
    return 1;
    }

    free(program_buffer);

    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

    kernel = clCreateKernel(program, KERNEL_FUNC, &err);

    if(err)
    {
    perror("Cannot create kernel");
    return 1;
    }

    queue = clCreateCommandQueue(context, device, CL_QUEU_PROFILING_ENABLE, &err);

    if(err)
    {
    perror("Cannot create command queue");
    return 1;
    }

    vec_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                sizeof(float) * ARRAY_SIZE, vec, &err);
    result_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float)*ARRAY_SIZE, NULL, &err);

    if(err)
    {
    perror("Cannot create the vector buffer");
    return 1;
    }

    clSetKernelArg(kernel, 0, sizeof(cl_mem), &vec_buffer);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &result_buffer);

    size_t global_size = ARRAY_SIZE;
    size_t local_size = 0;

    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, &prof_event);

    clEnqueueReadBuffer(queue, result_buffer, CL_TRUE, 0, sizeof(float)*ARRAY_SIZE, &result, 0, NULL, NULL);
    clFinish(queue);



     clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START,
           sizeof(time_start), &time_start, NULL);
     clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END,
           sizeof(time_end), &time_end, NULL);
     total_time += time_end - time_start;

    printf("\nAverage time in nanoseconds = %lu\n", total_time/ARRAY_SIZE);



    clReleaseMemObject(vec_buffer);
    clReleaseMemObject(result_buffer);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(queue);
    clReleaseProgram(program);
    clReleaseContext(context);

    free(vec);

    return 0;
}

而内核是：

__kernel void float_sum(__global float* vec,__global float* result){
    int gid = get_global_id(0);
    result[gid] = vec[gid] * vec[gid];
}

现在，结果如下：

秒数：0.010000 <- 这是 C 代码的时间

平均纳秒时间 = 140737284 <- OpenCL 函数

0.1407 秒是 OpenCL 时间内核执行的时间，比 C 函数更长，这正确吗？因为我认为 OpenCL 应该比 C 非并行算法更快...

我对这些结果感到非常惊讶，特别是因为您正在将OpenCL速度除以数组的大小。 您确定您正在正确计时代码吗？ 您使用的是Windows还是Linux？ 您使用的GPU是什么？ - KLee1
你是否正在尝试使用这个特定的例子？如果你使用float4类型，你可以在单个操作中进行点积和求和4个值。我假设你不是在寻找这样的优化，而是在使用OpenCL指针。还有一个平方和的公式可以使用。 - mfa
3个回答
3

在GPU上执行并行代码不一定比在CPU上执行更快。需要考虑的是，除了计算之外，你还需要将数据从GPU内存传输到CPU内存中，反过来也一样。

在您的例子中，您正在传输2 * N个项目并在并行执行一个O（N）操作，这是GPU非常低效的使用方式。因此，在这种特定的计算中，CPU很可能更快。

有没有办法使内核更有效率？我在考虑并行化数据... - user1319384
@facundo：不是很行。恐怕你唯一的选择就是在GPU上运行更重要的计算。例如，你可以尝试实现矩阵乘法。 - Tudor
1
这并没有什么意义。只有内核执行时间被测量了。它没有考虑数据传输。 - KLee1
一块普通的GPU在这个任务上往往会比高端CPU表现更出色。传输时间没有被测量，但只有4M的上传/下载速度。这个例子有点牵强。前N个平方数的和=（N （N + 1）（2N + 1））/ 6。 - mfa
2

为了帮助其他人，简单介绍如何使用OpenCL对内核运行时进行分析。

启用分析模式：

cmdQueue = clCreateCommandQueue(context, *devices, CL_QUEUE_PROFILING_ENABLE, &err);

内核剖析：

cl_event prof_event; 
clEnqueueNDRangeKernel(cmdQueue, kernel, 1 , 0, globalWorkSize, NULL, 0, NULL, &prof_event);

读取性能数据：

cl_ulong ev_start_time=(cl_ulong)0;     
cl_ulong ev_end_time=(cl_ulong)0;   

clFinish(cmdQueue);
err = clWaitForEvents(1, &prof_event);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &ev_start_time, NULL);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &ev_end_time, NULL);

计算内核执行时间：

float run_time_gpu = (float)(ev_end_time - ev_start_time)/1000; // in usec

你的方法与 

total_time/ARRAY_SIZE

这不是你想要的。它会给出每个工作项的运行时间。

操作/时间以纳秒为单位将给出GFLOPS（每秒十亿次浮点运算次数）。

2

你的应用程序存在一个大问题：

size_t global_size = ARRAY_SIZE;
size_t local_size = 0;

您正在创建单项工作组，这将导致大部分GPU处于闲置状态。在许多情况下，使用单项工作组只会利用 GPU 的 1/15。相反，请尝试以下方法：
size_t global_size = ARRAY_SIZE / 250; //important: local_size*global_size must equal ARRAY_SIZE
size_t local_size = 250; //a power of 2 works well. 250 is close enough, yes divisible by your 1M array size

现在您正在创建大型组，以更好地饱和图形硬件的ALU。目前您的内核可以正常运行，但是还有一些方法可以优化内核部分。
内核优化：将ARRAY_SIZE作为附加参数传递到内核中，并使用更少但更优化的组大小。这样还可以消除需要local_size*global_size恰好等于ARRAY_SIZE的需求。在此内核中，工作项的全局ID从未被使用过，也不需要，因为总大小已经传入了。
__kernel void float_sum(__global float* vec,__global float* result,__global int count){
  int lId = get_local_id(0);
  int lSize = get_local_size(0);
  int grId = get_group_id(0);
  int totalOps = count/get_num_groups(0);
  int startIndex = grId * totalOps;
  int maxIndex = startIndex+totalOps;
  if(grId == get_num_groups(0)-1){
    endIndex = count;
  }
  for(int i=startIndex+lId;i<endIndex;i+=lSize){
    result[i] = vec[i] * vec[i];
  }
}

现在你可能会认为对于这样一个简单的内核来说有太多的变量了。请记住，每次执行内核都会对数据进行多个操作，而不仅仅是一个。在我的Radeon 5870（20个计算单元）上使用以下值，每个工作项最终将在其for循环中计算781或782个值。每个组将计算50000个数据。我使用的变量的开销远小于创建4000个工作组或100万个工作项的开销。

size_t global_size = ARRAY_SIZE / numComputeUnits;
size_t local_size = 64; //also try other multiples of 16 or 64 for gpu; or multiples of your core-count for a cpu kernel

点击此处查看如何获取numComputeUnits的值

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接