为什么我的 OpenCL 内核在 nVidia driver 而不是 Intel 上失败（可能是 driver 错误）？

Question

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

很早就发现一个很奇怪的问题，如果我的内核大小正好是55，第27个内核线程将不会执行。此问题仅在使用nVidiadriver和GPU加速时出现。当我运行它在 CPU 上使用英特尔 driver 时，我发现第 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号）。

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

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 进行编程时是否可能存在一些我不知道的陷阱，或者我是否不幸发现了 driver 错误？

Hardware/OS 规格

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

Answer 1

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

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

Answer 2

这是 NVIDIA 的回复，一个是解决方法，一个是解决方案。我们只是 post 在我们的错误系统中，但没有得到你的回复，所以我们 post 在这里工作 around/solution。谢谢！

1.解决方法： 我们在本地重现了这个问题，并从我们的开发团队中提出了一个解决方法，请尝试下面的修改，让我们知道它是否有效。谢谢你的耐心。

您可以尝试以下解决方法。在文件 testIndex.cl 中将 sqrt() 改为使用 pow()，请参阅下面的代码片段。

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

在第 7.4 节中，table 说： sqrt <= 3 ulp

这里产生2个值：根 = 15.0000009537 根 = 15.0000000000 仅相差 1ULP，这是标准允许的。有关 ULP 的介绍，请参阅 https://en.wikipedia.org/wiki/Unit_in_the_last_place。您基本上需要在 opencl 编译时使用“-cl-fp32-correctly-rounded-divide-sqrt”选项，方法是指定 program.build(devices, "-cl-fp32-correctly-rounded-divide-sqrt");

为什么我的 OpenCL 内核在 nVidia driver 而不是 Intel 上失败（可能是 driver 错误）？

Why does my OpenCL kernel fail on the nVidia driver, but not Intel (possible driver bug)?

gpgpu

nvidia

opencl