使用 CUDA 4.2 和驱动程序 295.41 的非常有趣的行为

Posted 2023-04-15

【中文标题】使用 CUDA 4.2 和驱动程序 295.41 的非常有趣的行为

【英文标题】：very interesting behaviour using CUDA 4.2 and driver 295.41

【发布时间】：2012-08-20 21:38:51

【问题描述】：

在 Linux 上使用 CUDA 4.2 和驱动程序 295.41 时，我目睹了一个非常有趣的行为。 代码本身无非就是找到一个随机矩阵的最大值并将位置标记为1。

#include <stdio.h>
#include <stdlib.h>

const int MAX = 8;

static __global__ void position(int* d, int len) 
    int idx = threadIdx.x + blockIdx.x*blockDim.x;
    if (idx < len) 
        d[idx] = (d[idx] == MAX) ? 1 : 0;


int main(int argc, const char** argv) 
    int colNum = 16*512, rowNum = 1024;
    int len = rowNum * colNum;

    int* h = (int*)malloc(len*sizeof(int));
    int* d = NULL;
    cudaMalloc((void**)&d, len*sizeof(int));

    // get a random matrix
    for (int i = 0; i < len; i++) 
        h[i] = rand()%(MAX+1);
       

    // launch kernel
    int threads = 128;
    cudaMemcpy(d, h, len*sizeof(int), cudaMemcpyHostToDevice);
    position<<<(len-1)/threads+1, threads>>>(d, len);
    cudaMemcpy(h, d, len*sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(d);
    free(h);
    return 0;

当我设置 rowNum = 1024 时，代码根本不起作用，就好像内核从未启动过一样。 如果 rowNum = 1023，一切正常。

并且这个 rowNum 值在某种程度上与块大小（在本例中为 128）进行了卷积，如果我将块大小更改为 512，则行为发生在 rowNum = 4095 和 4096 之间。

我不太确定这是一个错误还是我错过了什么？

【参考方案1】：

您应该始终在调用 CUDA 函数后检查错误。例如，在您的代码中，invalid configuration argument 错误发生在内核启动期间。

这通常意味着网格或块的尺寸无效。

使用colNum = 16*512, rowNum = 1024，您正在尝试运行 65536 个块 x 128 个线程，超过了最大网格尺寸（对于具有计算能力 1.x 和 2.x 的 GPU，它是 65535 个块，不确定 3.x）。

如果您需要运行更多线程，您可以增加块大小（您已经尝试过并且效果不错）或使用 2D/3D 网格（3D 仅适用于计算能力为 2.0 或更高的设备）。

【讨论】：

是的，在这个特定的例子中你是对的，但是，正如我之前提到的，即使我设置 rowNum = 4096，块大小为 512，它也不起作用。

@user1624864 使用这样的配置（colNum = 16*512，rowNum = 4096，threads = 512）启动配置是65536 x 512，它也不应该工作

好的，我明白你的意思了。谢谢，大网格尺寸是该行为的原因。但是，我只是在帖子中省略了所有错误检查代码，而不是在生产中。那么问题就变成了一个普遍的问题，SDK提供的“cuda安全调用”方法以及驱动错误报告是否可靠？

虽然我有错误检查模板，但之前没有关于此错误的投诉。

@user1624864：跑得更快 = 启动失败并且根本没有运行。您显然没有正确执行错误检查。内核启动后直接调用cudaPeekAtLastError 并检查返回状态。这将检测到无效的启动配置。