运行时OpenCL错误计算矩阵乘法

Question

过去几天我一直在调试，无法运行OpenCL矩阵乘法内核。每当我运行该程序时，来自GPU的输出都会产生类似于-198746573.0000的大负数。我想知道是否有HPC经验的人可以在我的代码中指出错误，或者是否与驱动程序有关。

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

#define widthA 2 
#define heightA 2 

#define widthB heightA 
#define heightB 2 

#define widthC widthA 
#define heightC heightB 

#ifdef __APPLE__ 
#include <OpenCL/opencl.h> 
#else 
#include <opencl.h> 
#endif 

#define MEM_SIZE (128) 
#define MAX_SOURCE_SIZE (0x100000) 

int main() 
{ 
    float * A = (float *)malloc(sizeof(float)*widthA*heightA); 
    float * B = (float *)malloc(sizeof(float)*widthB*heightB); 
    float * C = (float *)malloc(sizeof(float)*widthC*heightC); 
    float * Res = (float *)malloc(sizeof(float)*widthC*heightC); 
    float * D= (float *)malloc(sizeof(float)*widthC*heightC); 

    float ref[widthC][heightC]; 

    int i, j, k; 

    FILE * fp1 = fopen("matAdata.txt", "w"); 
    if (!fp1) { 
    fprintf(stderr, "Failed to open matAdata.\n"); 
    exit(1); 
    } 

    for(i = 0;i < widthA; i++) 
    { 
     for(j=0;j < heightA; j++)  { 
      float p=(rand()%100)/7.0; 
      //*(A+i*heightA+j)=rand()%100 + p; 
      *(A+i*heightA+j)=4.0; 
      fprintf(fp1, "%f ",*(A+i*heightA+j)); 
     } 
     fprintf(fp1, "\n"); 
    } 
    fclose(fp1); 

    fp1 = fopen("matBdata.txt", "w"); 
    if (!fp1) { 
    fprintf(stderr, "Failed to open matAdata.\n"); 
    exit(1); 
    } 


    for(i = 0;i < widthB; i++) 
    { 
     for(j=0; j < heightB; j++)  { 
      float p=(rand()%100)/7.0; 
      //*((B+i*heightB+j))=rand()%100 + p; 
      *((B+i*heightB+j))=4.0; 
      fprintf(fp1, "%f ",*(B+i*heightA+j)); 
     } 
     fprintf(fp1, "\n"); 
    } 
    fclose(fp1); 

    cl_device_id device_id = NULL; 
    cl_context context = NULL; 
    cl_command_queue command_queue = NULL; 
    cl_mem memobjA = NULL; 
    cl_mem memobjB = NULL; 
    cl_mem memobjC = NULL; 
    cl_mem rowA = NULL; 
    cl_mem colC = NULL; 
    cl_program program = NULL; 
    cl_kernel kernel = NULL; 
    cl_platform_id platform_id[10]; 
    cl_platform_id platform = NULL; 
    cl_uint ret_num_devices; 
    cl_uint ret_num_platforms; 
    cl_int ret; 
    cl_event GPUDone[0]; 
    //char string[MEM_SIZE]; 

    FILE *fp; 
    char fileName[] = "matrixMultiplication.cl"; 
    char *source_str; 
    size_t source_size; 
    int row = widthA; 
    int col = heightC; 
    /* Load the source code containing the kernel*/ 
    fp = fopen(fileName, "r"); 
    if (!fp) { 
    fprintf(stderr, "Failed to load kernel.\n"); 
    exit(1); 
    } 
    source_str = (char*)malloc(MAX_SOURCE_SIZE); 
    source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp); 
    fclose(fp); 

    /* Get Platform and Device Info */ 
    ret = clGetPlatformIDs(10, platform_id, &ret_num_platforms); 

    char cBuffer[1024]; 
    cl_uint c; 

    for(c = 0; c < ret_num_platforms; c++) 
    { 
    clGetPlatformInfo(platform_id[c], CL_PLATFORM_NAME, 1024, &cBuffer, NULL); 
    if (strstr(cBuffer, "NVIDIA") != NULL) 
    { 
     platform = platform_id[c]; 
     break; 
    } 

    } 

    printf("Found Platform %s\n", cBuffer); 

    ret = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device_id, &ret_num_devices); 

    printf("Found %d devices.\n", ret_num_devices); 

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

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

    /* Create Memory Buffer */ 
    memobjA = clCreateBuffer(context, CL_MEM_READ_ONLY, widthA * heightA * sizeof(float), NULL, &ret); 
    memobjB = clCreateBuffer(context, CL_MEM_READ_ONLY, widthB * heightB * sizeof(float), NULL, &ret); 
    memobjC = clCreateBuffer(context, CL_MEM_READ_WRITE, widthC * heightC * sizeof(float), NULL, &ret); 
    rowA = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int), NULL, &ret); 
    colC = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int), NULL, &ret); 

    // Copy the lists A and B to their respective memory buffers 
    ret = clEnqueueWriteBuffer(command_queue,memobjA, CL_TRUE, 0, 
      widthA * heightA * sizeof(float), A, 0, NULL, NULL); 
    ret = clEnqueueWriteBuffer(command_queue, memobjB, CL_TRUE, 0, 
      widthB * heightB * sizeof(float), B, 0, NULL, NULL); 
    ret = clEnqueueWriteBuffer(command_queue, rowA, CL_TRUE, 0, sizeof(int), &row, 0, NULL, NULL); 
    ret = clEnqueueWriteBuffer(command_queue, colC, CL_TRUE, 0, sizeof(int), &col, 0, NULL, NULL); 

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

    /* Build Kernel Program */ 
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); 

    /* Create OpenCL Kernel */ 
    kernel = clCreateKernel(program, "matrixMultiplication", &ret); 

    /* Set OpenCL Kernel Arguments */ 
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjA); 
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjB); 
    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&memobjC); 
    ret = clSetKernelArg(kernel, 3, sizeof(int), (void *)&row); 
    ret = clSetKernelArg(kernel, 4, sizeof(int), (void *)&col); 
    /* Execute OpenCL Kernel */ 

    //ret = clEnqueueTask(command_queue, kernel, 0, NULL,NULL); 
    size_t globalThreads[2] = {widthA, heightB}; 
    size_t localThreads[2] = {16,16}; 

    clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, globalThreads, localThreads, 0, NULL, NULL); 
    //clFlush(command_queue); 
    //clFinish(command_queue); 

    /* Copy results from the memory buffer */ 
    ret = clEnqueueReadBuffer(command_queue, memobjC, CL_TRUE, 0, 
          widthA * heightC * sizeof(float), Res, 0, NULL, &GPUDone[0]); 

    printf("Buffer Read ended with %d.\n", ret); 
    clWaitForEvents(1, GPUDone); 

    fp1 = fopen("matGPURes.txt", "w"); 
    if (!fp1) { 
    fprintf(stderr, "Failed to open matAdata.\n"); 
    exit(1); 
    } 

    printf("\nResult\n"); 
    for(i = 0;i < widthA; i++) 
    { 
     for(j=0;j < heightC; j++) 
     { 

      fprintf(fp1, "%f ",*(Res+i*heightC+j)); 
      ref[i][j] = *(Res+i*heightC+j); 
      printf("GPU Output: %f\n", *(Res+i*heightC+j)); 
     } 
     fprintf(fp1, "\n"); 
    } 
    fclose(fp1); 

    ret = clFlush(command_queue); 
    ret = clFinish(command_queue); 
    ret = clReleaseKernel(kernel); 
    ret = clReleaseProgram(program); 
    ret = clReleaseMemObject(memobjA); 
    ret = clReleaseMemObject(memobjB); 
    ret = clReleaseMemObject(memobjC); 
    ret = clReleaseCommandQueue(command_queue); 
    ret = clReleaseContext(context); 
    ret = clReleaseEvent(GPUDone[0]); 

    free(source_str); 

    float sum=0.0; 

    for(i = 0;i < widthA; i++) 
    { 
     for(j = 0; j < heightC; j++) 
     { 
      sum = 0; 
      for(k = 0; k < widthB; k++) 
      { 
       sum += A[i*col+k] * B[k*row+j]; 
       printf("Multiplying A: %f, B: %f\n", A[i*col+k], B[k*row+j]); 
      } 
     D[i*heightC+j] = sum; 
     } 

    } 

    fp1 = fopen("matNormalMultiplicationRes.txt", "w"); 

    if (!fp1) { 
    fprintf(stderr, "Failed to open matNormalMultiplicationRes.txt\n"); 
    exit(1); 
    } 

    for(i = 0; i<widthA; i++) 
    { 
     for(j = 0; j<heightA; j++) 
     { 
      if (ref[i][j] != D[i*heightA+j]) 
      { 
       printf("Calculation error[ CPU: %f, GPU: %f ]\n", D[i*heightA+j], ref[i][j]); 
      } 
     } 
    } 

    printf("\nResult\n"); 
    for(i = 0;i < widthA; i++) 
    { 
     for(j=0;j < heightC; j++) 
     { 
      fprintf(fp1, "%f ",*(D+i*heightC+j)); 

     } 
     fprintf(fp1, "\n"); 
    } 
    free(A); 
    free(B); 
    free(C); 
    free(D); 
    free(Res); 
    return 0; 
} 

这里是内核

#define BLOCK_SIZE 16 

__kernel 
void matrixMultiplication(__global float* A, __global float* B, __global float* C, int wA, int wB) 
{ 
    //int i = get_global_id(0); 
    //int j = get_global_id(1); 

    float Csub = 0.0f;   

    int bx = get_group_id(0); 
    int by = get_group_id(1); 

    int tx = get_local_id(0); 
    int ty = get_local_id(1); 

    int aBegin = wA * BLOCK_SIZE * by; 
    int aEnd = aBegin + wA - 1; 
    int aStep = BLOCK_SIZE; 

    int bBegin = BLOCK_SIZE * bx; 
    int bStep = BLOCK_SIZE * wB; 

    for (int a = aBegin, b=bBegin; 
     a <= aEnd; 
     a += aStep, b+=bStep) 
    { 
     __local float As[BLOCK_SIZE][BLOCK_SIZE]; 
     __local float Bs[BLOCK_SIZE][BLOCK_SIZE]; 

     As[ty][tx] = A[a + wA * ty + tx]; 
     Bs[ty][tx] = B[b + wB * ty + tx]; 
     barrier(CLK_LOCAL_MEM_FENCE); 

     for(int k = 0; k < BLOCK_SIZE; ++k) 
      Csub += As[ty][k] * Bs[k][tx]; 
     barrier(CLK_LOCAL_MEM_FENCE); 

    } 

    int c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx; 
    C[c + wB * ty + tx] = Csub; 
    /* 
    float value=0; 
    for (int k = 0; k < widthA; k++) 
    { 
     value = value + A[k + j * widthA] * B[k*widthB + i]; 
    } 
    C[i + widthA * j] = value; 
    */ 
} 

我有双重检查了一遍又一遍，但根本无法找到任何错误。我想在确定它的驱动程序问题之前确保它没有代码错误。

谢谢！

Answer 1

你真的需要这样一个复杂的内核吗？如果你真的想做简单的矩阵乘法运算 ，你可以编写一个简单的内核，这很容易调试。

__kernel void matrixMultiplication (__global float* A, 
             __global float* B, 
             __global float* C, 

             int widthA, int widthB) 
{ 
    //y direction 
    int row = get_global_id(1); 

    int col = get_global_id(0); 

    float cSum = 0.0f; 

    //calculate the result 
    for (int i=0; i<widthA; i++) 
    { 
     cSum += A[row*widthA+ i] * B[i*widthB+col]; 
    } 

    C[row*widthB+col] = cSum; 
} 

Answer 2

检查您的主机的功能。这里有几件事让你开始...

1）你不需要创建一个缓冲区并将其排入标量常量int，如行和列。只需将其设置为内核参数。

2）等待带有事件的clEnqueueNDRangeKernel。你想确定calc已经完成。

3）在内核中添加一个printf语句来打印选定的值以查看输入和输出值是否符合您的期望值。

尝试

如果（get_local_id（0）％8 == 0）{

printf some useful value of a,b,c 

}

3）尝试用哑内核拷贝的输入阵列的主机代码到输出数组。这将确认你有缓冲区创建处理和条件读/写代码正确！

Answer 3

此案例可能已关闭，但为了google-comers： 不应该在主机上显式声明共享内存并将其作为内核参数传递给源代码？在这种情况下，__local关键字不是您正在寻找的关键字。

有关详细说明，请参阅How to declare local memory in OpenCL?的文章。