opencl - OpenCL 内核的随机 NaN 和错误结果

Question

我正在尝试实现一种通用矩阵-矩阵乘法 OpenCL 内核，该内核符合 C = α*A*B + β*C。

内核

我在网上做了一些研究，并决定使用 this website 中的修改内核。作为起点。我所做的主要修改是本地内存作为工作空间的分配现在是动态的。下面是我写的内核:

__kernel
void clkernel_gemm(const uint M, const uint N, const uint K, const float alpha,
                   __global const float* A, __global const float* B, const float beta, 
                   __global float* C, __local float* Asub, __local float* Bsub) {

  const uint row = get_local_id(0);
  const uint col = get_local_id(1);
  const uint TS = get_local_size(0); // Tile size
  const uint globalRow = TS * get_group_id(0) + row; // Row ID of C (0..M)
  const uint globalCol = TS * get_group_id(1) + col; // Row ID of C (0..N)

  // Initialise the accumulation register
  float acc = 0.0f;

  // Loop over all tiles
  const int numtiles = K / TS;
  for (int t = 0; t < numtiles; t++) {
    const int tiledRow = TS * t + row;
    const int tiledCol = TS * t + col;
    Asub[col * TS + row] = A[tiledCol * M + globalRow];
    Bsub[col * TS + row] = B[globalCol * K + tiledRow];

    barrier(CLK_LOCAL_MEM_FENCE);

    for(int k = 0; k < TS; k++) {
      acc += Asub[k * TS + row] * Bsub[col * TS + k] * alpha;
    }

    barrier(CLK_LOCAL_MEM_FENCE);
  }

  C[globalCol * M + globalRow] = fma(beta, C[globalCol * M + globalRow], acc);
}

图 block 大小 (TS) 现在是调用代码中定义的值，如下所示:

  // A, B and C are 2D matrices, their cl::Buffers have already been set up
  // and values appropriately set.

  kernel.setArg(0, (cl_int)nrowA);
  kernel.setArg(1, (cl_int)ncolB);
  kernel.setArg(2, (cl_int)ncolA);
  kernel.setArg(3, alpha);
  kernel.setArg(4, A_buffer);
  kernel.setArg(5, B_buffer);
  kernel.setArg(6, beta);
  kernel.setArg(7, C_buffer);
  kernel.setArg(8, cl::Local(sizeof(float) * nrowA * ncolB));
  kernel.setArg(9, cl::Local(sizeof(float) * nrowA * ncolB));

  cl::NDRange global(nrowA, ncolB);
  cl::NDRange local(nrowA, ncolB);

  status = cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), global, local);

问题

我遇到的问题是，我编写的单元测试(用 Google 的 gtest 编写)会随机失败，但仅限于这个特定的内核。 (我在同一个 .cl 源文件中有 20 个其他内核，它们 100% 通过测试)

我有一个测试，将 1x4 浮点矩阵 {0.0, 1.0, 2.0, 3.0} 与其转置版本 {{0.0}, {1.0}, {2.0} ，{3.0}}。预期输出为 {14.0}。

但是，我可能只有 75% 的时间才能得到正确的结果。

有时，我可以获得 23.0 (GTX 970)、17.01 (GTX 750) 或只是 -nan 和 0.0(所有 3 台设备)。奇怪的是，各自的错误结果似乎是各设备所独有的；例如，我似乎无法在 Intel CPU 或 GTX 750 上获得 23.0。

我很困惑，因为如果我犯了算法或数学错误，那么这个错误应该是一致的；相反，我只是随机得到不正确的结果。

我在这里做错了什么？

我尝试过的事情

• 我已经验证进入内核的数据是正确的。
• 我尝试将 __local 内存初始化为 0.0，但这会导致所有结果变得错误(但坦率地说，我不太确定如何正确初始化它)
• 我编写了一个测试程序，该程序仅执行该内核，以排除与程序其余部分交互的任何竞争条件，但错误仍然发生。

其他注意事项

• 我正在使用直接从 the Github page 检索的 C++ 包装器.
• 为了使用包装器，我定义了 CL_HPP_MINIMUM_OPENCL_VERSION 120 和 CL_HPP_TARGET_OPENCL_VERSION 120。
• 我正在使用 -cl-std=CL1.2 标志编译内核。
• 所有 cl::Buffer 均仅使用 CL_MEM_READ_WRITE 标志创建。
• 我正在 Ubuntu 16.04、Ubuntu 14.04 和 Debian 8 上测试此功能。
• 我已在 Intel CPU 上对此进行了测试，编号为 Intel OpenCL Runtime 16.1 for Ubuntu安装。运行时报告它最高支持 OpenCL 1.2
• 我已在 Nvidia GTX 760 和 970 上对此进行了测试。Nvidia 仅支持 OpenCL 1.2。
• 所有 3 个平台都出现相同的问题，但出现的频率各不相同。

Answer 1

这看起来很复杂。有几件事需要解决，它们不适合评论，所以我会将所有这些作为答案发布，即使它还没有解决您的问题。

---

I am baffled because if I have made an algorithmic or mathematical mistake, the mistake should be consistent; instead I am getting incorrect results only randomly.

这种行为是竞争条件的典型指标。

---

I have tried to initialize both __local memory to 0.0, but this causes all results to become wrong (but frankly, I'm not really sure how to initialize it properly)

其实这是一件好事。最后我们有了一定的一致性。

---

初始化本地内存

可以使用工作项来初始化本地内存，例如如果您有一个包含 16 个项目的一维工作组，并且您的本地内存由 16 个浮点组成，只需执行以下操作:

local float* ptr = ...          // your pointer to local memory
int idx = get_local_id(0);      // get the index for the current work-item
ptr[idx] = 0.f;                 // init with value 0
barrier(CLK_LOCAL_MEM_FENCE);   // synchronize local memory access within workgroup

如果您的本地内存较大，例如64 个 float ，您将不得不使用一个循环，其中每个工作项初始化 4 个值，至少这是最有效的方法。但是，没有人会阻止您使用每个工作项来初始化本地内存中的每个值，即使这完全是无稽之谈，因为您实际上多次初始化它。

---

您的更改

original algorithm看起来它是专门为使用二次瓷砖而设计的。

__local float Asub[TS][TS];
__local float Bsub[TS][TS];

不仅如此，本地内存的大小也与工作组大小相匹配，在他们的示例中为 32x32。 当我查看本地内存的内核参数时，我可以看到您使用了原始算法中定义为 M 和 N 的参数。这似乎不正确。

更新 1

由于您没有描述原始算法是否适合您，因此您应该执行以下操作来查找错误:

• 创建一组测试数据。确保仅使用原始算法实际支持的数据大小(例如最小大小、x 的倍数等)。此外，请使用大型数据集，因为某些错误仅在调度多个工作组时才会显示。
• 将原始的、未更改的算法与您的测试数据集结合使用并验证结果。
• 仅更改算法，使用动态本地内存大小代替固定大小本地内存，但确保其大小与固定大小方法相同。这是您尝试过的方法，但我认为由于我在“您的更改”下描述的内容，它失败了。