为什么我的 OpenCL 内核在 nVidia 驱动程序上失败，而不是 Intel（可能的驱动程序错误）？答案

【问题标题】：Why does my OpenCL kernel fail on the nVidia driver, but not Intel (possible driver bug)?为什么我的 OpenCL 内核在 nVidia 驱动程序上失败，而不是 Intel（可能的驱动程序错误）？
【发布时间】：2018-03-06 15:54:35
【问题描述】：

我最初编写了一个 OpenCL 程序来计算非常大的厄米矩阵，其中内核计算矩阵中的一对条目（上三角部分和下三角补码）。

很早的时候，我发现了一个非常奇怪的问题，如果我的内核大小正好是 55，那么第 27 个内核线程将不会执行。此问题仅在使用 nVidia 驱动程序和 GPU 加速时发生。当我在 CPU 上使用 Intel 驱动程序运行它时，我发现第 27 个内核线程执行得很好。越来越小的内核大小似乎没有表现出问题。

认为这可能是我的代码中的某些东西，我将我的问题提炼为以下非常简单的内核：

__kernel void testIndex(__global float* outMatrix, unsigned int sizeN)
{
    //k is the linear kernel ID (related to but not exactly the linear index into the outMatrix)
    int k = get_global_id(0);
    //i'th index (Row or Y)
    int i = floor((2 * sizeN+1 - sqrt((float)((2 * sizeN + 1) * (2 * sizeN + 1) -8 * k) )) /2);

    //j'th index (Column or X)
    int j = k - sizeN * i + i * (i - 1) / 2;
    j += i;

    //Index bounds check... If we're greater than sizeN, we're an idle core.
    //(OpenCL will queue up a fixed block size of worker threads, some of them may be out of bounds)
    if(j >= sizeN || i >= sizeN)
    {
        return;
    }

    //Identity case. The original kernel did some special stuff here,
    //but I've just replaced it with the K index code.
    if(i == j)
    {
        outMatrix[i * sizeN +j] = k;
        return;
    }

    outMatrix[i * sizeN + j] = k;

    //Since we only have to calculate the upper triangle of our matrix,
    //(the lower triangle is just the complement of the upper),
    //this test sets the lower triangle to -9999 so it's easier to see
    //how the indexing plays out...

    outMatrix[j * sizeN + i] = -9999.0;

 }

outMatrix 是输出矩阵，sizeN 是方阵在一边的大小（即矩阵是 sizeN x sizeN）。

我使用以下主机代码计算并执行我的内核大小：

size_t kernelSize = elems * (elems + 1) / 2;
cl::NDRange globalRange(kernelSize);
cl::NDRange localRange(1);
cl::Event event;

clCommandQueue.enqueueNDRangeKernel(testKernel, cl::NullRange, globalRange, cl::NullRange, NULL, &event);
event.wait();

elems 与 sizeN 相同（即矩阵大小的平方根）。在这种情况下，elems = 10（因此内核大小为 55）。

如果我打印出我读回的矩阵，我会得到以下信息（使用 boost ublas 矩阵格式）：

[10,10] ((    0,     1,     2,     3,     4,     5,     6,     7,     8,    9),
        ((-9999,    10,    11,    12,    13,    14,    15,    16,    17,   18),
        ((-9999, -9999,    19,    20,    21,    22,    23,    24,    25,   26),
        ((-9999, -9999, -9999,  JUNK,    28,    29,    30,    31,    32,   33),
        ((-9999, -9999, -9999, -9999,    34,    35,    36,    37,    38,   39),
        ((-9999, -9999, -9999, -9999, -9999,    40,    41,    42,    43,   44), 
        ((-9999, -9999, -9999, -9999, -9999, -9999,    45,    46,    47,   48),
        ((-9999, -9999, -9999, -9999, -9999, -9999, -9999,    49,    50,   51),    
        ((-9999, -9999, -9999, -9999, -9999, -9999, -9999, -9999,    52,   53),   
        ((-9999, -9999, -9999, -9999, -9999, -9999, -9999, -9999, -9999,   54))

其中“JUNK”是一个随机值，基于当时内存中发生的任何事情。这当然是可疑的，因为 27 基本上是内核的中间点。

为了完整起见，使用以下代码回读矩阵结果：

boost::scoped_array<float> outMatrixReadback(new float[elems * elems]);
clCommandQueue.enqueueReadBuffer(clOutputMatrixBuffer, CL_TRUE, 0, elems * elems * sizeof(float), outMatrixReadback.get());

我做出（也许是不正确的）假设，即由于代码在 Intel CPU 上执行良好，代码本身不存在一些基本错误。

那么，在 nVidia 卡上编程 OpenCL 时可能有一些我不知道的问题，还是我很不幸发现了驱动程序错误？

硬件/操作系统规格

nVidia GTX 770
RHEL 服务器版本 6.4（圣地亚哥）
英特尔 OpenCL 1.2 4.4.4.0.134 SDK 标头
nVidia GeForce 驱动程序 384.69
英特尔至强 CPU E6520 @ 2.4 GHz

【问题讨论】：

CPU 和 GPU 之间存在许多差异 - 主要是内存一致性和算术精度。这里没有明显的罪魁祸首，但是 - 疯狂的建议：这可能是平方根的精度问题吗？将floor() 更改为round() 有什么不同吗？ GPU 上的平方根是出了名的不精确，但我确实希望 sqrt(225.0) = 15.0 成立……
据我所知，对于 k=27，内核根本没有执行。如果我做一些像“if(k == 27) outMatrix[0] = 4242; return;”这样简单的事情，我没有得到任何迹象表明 K 曾经等于 27。
如果您注释掉内核中除if (k == 27) printf("Thread reached\n"); 之外的所有其他代码。输出是什么？
奇怪的是，that 可以工作，但是如果你有任何内核线程做任何其他操作，27 将无法执行。
您是否尝试过明确地将“自然”数量的工作项排入队列？例如。 64 而不是 55。奇怪的工作规模可能是实现中测试不佳的代码路径。

标签： opencl nvidia gpgpu

【解决方案1】：

在与 nVidia 讨论后，技术代表确认这既是可重复的，也是驱动程序错误。已提交错误报告，但不幸的是，我被告知 nVidia 没有专门的 OpenCL 开发团队，因此无法提供修复时间表。

编辑： 在最终收到 nVidia 的回复后，解决方法显然是在 CL 内核中使用 pow() 而不是 sqrt()，因为 sqrt() 显然是错误的来源。

【讨论】：

【解决方案2】：

这里是 NVIDIA 的回复，一个正在解决，一个是解决方案。我们只是在我们的错误系统中发布，但没有得到您的回复，所以我们在此处发布解决方案/解决方案。谢谢！

1.解决方法： 我们在此问题上进行了本地复制，并从我们的开发团队中提出了一种解决方法，请尝试以下修改并让我们知道它是否有效。感谢您的耐心。

您可以尝试以下解决方法。在文件 testIndex.cl 中将 sqrt() 改为使用 pow()，参见下面的 sn-p。

//i'th index (Row or Y)
// original version
// FAIL float sqrt
//int i = floor((2 * sizeN+1 - sqrt((float)((2 * sizeN + 1) * (2 * sizeN + 1) -8 * k) )) /2);
// PASS float pow
//int i = floor((2 * sizeN+1 - pow((float)((2 * sizeN + 1) * (2 * sizeN + 1) -8 * k), 0.5f)) /2);

2。解决方案：今天有新的问题解决方案，请查看以下描述和解决问题的方法，让我们知道它是否有效。谢谢。

OpenCL 1.2 规范，第 5.6.4.2 节说： -cl-fp32-正确舍入除法-sqrt。 clBuildProgram 或 clCompileProgram 的 -cl-fp32-correctly-rounded-divide-sqrt 构建选项允许应用程序指定程序源中使用的单精度浮点除法（x/y 和 1/x）和 sqrt 正确舍入.如果未指定此构建选项，则单精度浮点除法和 sqrt 的最小数值精度在 OpenCL 规范的第 7.4 节中定义。

在第 7.4 节中，该表显示： sqrt

这两个值在这里产生：根 = 15.0000009537 根 = 15.0000000000 仅在标准允许的 1ULP 上有所不同。有关 ULP 的介绍，请参阅 https://en.wikipedia.org/wiki/Unit_in_the_last_place。通过指定 program.build(devices, "-cl-fp32-correctly-rounded-divide-sqrt");

【讨论】：