为什么我的OpenCL内核在nVidia驱动程序上失败,而不是英特尔(可能的驱动程序错误)？

Question

我最初写了一个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时不知道,或者我不幸发现了驱动程序错误？

硬件/ OS规格

• nVidia GTX 770

• RHEL Server版本6.4(圣地亚哥)

• 英特尔OpenCL 1.2 4.4.4.0.134 SDK标头

• nVidia GeForce驱动程序384.69

• Intel Xeon CPU E6520 @ 2.4 GHz

Answer 1

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

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