为啥我的 OpenCL 内核在 nVidia 驱动程序上失败,而不是 Intel(可能的驱动程序错误)?

Posted

技术标签:

【中文标题】为啥我的 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();

elemssizeN 相同(即矩阵大小的平方根)。在这种情况下,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。奇怪的工作规模可能是实现中测试不佳的代码路径。 【参考方案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");

【讨论】:

以上是关于为啥我的 OpenCL 内核在 nVidia 驱动程序上失败,而不是 Intel(可能的驱动程序错误)?的主要内容,如果未能解决你的问题,请参考以下文章

使用 Nvidia 显卡安装 AMD OpenCL CPU 驱动程序

NVIDIA OpenCL 设备版本

如何在 Nvidia GPU 上调试 OpenCL?

从 OpenCL 内核修改 VBO 数据

OpenCL AMD 与 NVIDIA 性能对比

以前听老人说在n卡上想跑起来opencl程序必须先安装cuda,我没装cuda只安装了nvidia的显卡驱动,结果配置