c - 当矩阵大小变得太大时，用 OpenCL 编写的矩阵乘法内核不起作用

Question

我编写了一个 OpenCL 矩阵乘法内核，它将执行两个方阵的乘法运算。内核代码为

void kernel product(global const float* A, global const float* B, global float* C, int n){
  size_t kx=get_global_id(0);
  size_t ky=get_global_id(1);
  for(int i=0; i<n; i++){
     C[n*kx+ky]=C[n*kx+ky]+A[n*kx+i]*B[n*i+ky];
  } 
}

启动内核的主机代码是

  // create buffer on the context
  int n=1000;
  cl::Buffer buffer_A(context,CL_MEM_READ_ONLY,sizeof(float)*(n*n));
  cl::Buffer buffer_B(context,CL_MEM_READ_ONLY,sizeof(float)*(n*n));
  cl::Buffer buffer_C(context,CL_MEM_READ_WRITE,sizeof(float)*(n*n));

  float* A=new float[n*n];
  float* B=new float[n*n];
  float* C=new float[n*n];

  for (int i=0; i<n; i++) {
    for (int j=0; j<n; j++) {
      A[n*i+j]=2.0;
      B[n*i+j]=2.0;
    }
  }

  //create the kernel, and set the buffer argument  
  cl::Kernel kernel(program,"product");
  kernel.setArg(0, buffer_A);
  kernel.setArg(1, buffer_B);
  kernel.setArg(2, buffer_C);
  kernel.setArg(3, n);

  //build the queue
  cl::Device device_use=all_devices[0];
  cl::CommandQueue queue(context,device_use);

  // queue manipulation: step 1: write the input buffer
  queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(float)*(n*n), A);
  queue.finish();
  queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, sizeof(float)*(n*n), B);
  queue.finish(); 
  // queue manipulation: Step 2 run kernel
  queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(n,n), cl::NullRange);
  queue.finish();

请注意，A、B、C 是维度为 n*n 的方阵。我尝试在 Macbook pro 上的 Intel Iris 显卡上运行这个内核。当 n 较小时效果很好。但是，当 n 为 2000 或更大时，它会给出错误的结果。此 gpu 的最大全局工作大小为 (512,512,512)。所以2000*2000肯定不会超过最大值。当我尝试在 cpu 上运行内核时，无论 n 有多大，我总能得到正确的结果。所以内核应该是对的。对发生的事情有什么想法吗？

Answer 1

看来这里有几个问题。我会尝试解决所有这些问题(有些可能已经在我的评论中解决)。

OpenCL 不保证全局内存的正确初始化。有些设备可能会初始化为零，但有些不会。但是，您的代码确实依赖于此，因为您是在写入单个值之前从全局内存中读取的: C[n*kx+ky]=C[n*kx+ky]+A[n* kx+i]*B[n*i+ky];.此外，您还不必要地访问全局内存。您不应将中间结果保存在全局内存中，而应保存在快速私有(private)内存中(请参阅改进的内核代码，它也处理 C 未初始化的事实)。

您似乎不太清楚 OpenCL 本地和全局工作大小是如何处理的，所以我稍微谈谈这个。

工作规模限制(您的工作规模必须满足所有这些要求):

• CL_DEVICE_MAX_WORK_ITEM_SIZES 返回每个维度的最大本地工作大小。因此，您本地工作大小的每个维度都必须等于或小于相应的值。示例:CL_DEVICE_MAX_WORK_ITEM_SIZES 返回 [512,512,512]，因此 [512,2,1] 的本地工作大小是合法的，[2,512,1] 也是如此。但是 [1024,1,1] 是非法的，因为它违反了第一个维度的最大大小。

• CL_DEVICE_MAX_WORK_GROUP_SIZE 返回您的设备支持的每个工作组的最大工作项数，即您本地工作大小内的最大工作项数。如果 CL_DEVICE_MAX_WORK_GROUP_SIZE 返回 1024，则 [512,2,1] 是合法的，[1024,1,1] 也是合法的，但 [1024,2,1] 是非法的，因为 1024*2 > 1024。

• CL_KERNEL_WORK_GROUP_SIZE 返回您的设备为此特定内核支持的每个工作组的最大工作项数。这通常与 CL_DEVICE_MAX_WORK_GROUP_SIZE 相同，但对于使用大量私有(private)和/或本地内存的内核，它可能会更低。

• 您的全局工作规模必须是本地工作规模的倍数。如果矩阵的大小是 [2000,2000]，这似乎是一件微不足道的事情。您选择全局大小相同，OpenCL 会为您计算本地工作大小。我可能会是 [16,16]，因为它们是 2000 年的最大除数，并且仍然会产生低于 512 的局部工作大小。但请考虑一下:您的矩阵大小为 [905,905]。 OpenCL 将不得不选择 [1,1] 的本地工作大小，这在性能方面是有史以来最糟糕的情况(除非您的设备足够智能以弥补这种糟糕的工作大小)。 905 不能除以 1 以外的任何整数。请注意，我对此可能是错误的，但在阅读了很多关于 OpenCL 的内容之后，我怀疑这就是它“必须”计算工作大小的方式。因此，为了获得高性能，工作组通常不应小于 64，但在现代设备上 256 是一个非常好的值。因此，您应该根据这些值计算全局工作大小并调整您的内核，以便它可以处理比需要处理的元素更多的工作项。示例:您想要一个大小为 [16,16] = 256 的工作组，但您的矩阵有 1000 行和列。因此你的全局工作大小应该是 [1024,1024] 并且你的内核应该丢弃所有不需要的工作项。如果您仍希望 OpenCL 选择本地工作大小，只需将全局工作大小更改为 128 或 256 的倍数，以避免退化本地工作组大小。

内核代码:

void kernel product(global const float* A, global const float* B, global float* C, int n)
{
    size_t kx=get_global_id(0);
    size_t ky=get_global_id(1);

    // Discard work-items that are not needed.
    if(kx >= n || ky >= n)
        return;

    float result = 0.f;

    int idxC = n*kx+ky;

    for(int i=0; i<n; ++i)
    {
        int idxA = n*kx+i;
        int idxB = n*i+ky;

        result += A[idxA]*B[idxB];
    } 

    C[idxC] = result;
}

内核代码结束