超时后启动在 cudaDeviceSynchronize 中终止

Question

我尝试使用 3 维网格运行一个简单的程序，但由于某种原因，当我使用 cuda-memcheck 启动它时，它卡住了，超时后它终止了。该问题与短暂的超时无关，因为我只是将这种方式更改为 60 秒。

我运行的代码有一个 45x1575x1575 的网格，它运行一个空的 __global__ 函数。我的计算能力是 2.1，我使用标志 -maxrregcount=24 进行编译以限制设备功能可以使用的寄存器数量（在我的一些其他程序中看到，它使用占用计算器提供了最佳结果）

这是我的代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void stam(int a){

}

int main()
{


    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaError_t cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        return;
    }

    dim3 gridSize(45,1575,1575);
    stam<<<gridSize,224>>>(4);
    cudaStatus = cudaDeviceSynchronize(); // This function gets stuck
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!!");
        return;
    }

    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

最大网格尺寸不是 65535x65535x65535 吗？这里有什么问题？

编辑：只有当我使用 -G 标志编译它时它才会崩溃。否则只是慢，但不会超过 60 秒。

Answer 1

我刚刚在 cuda-memcheck 和 cuda-gdb 下的 C2050（功能 2.0）上运行您的代码，没有任何问题。这表明问题更可能与您的卡/设置有关，而不是代码本身。

如果您超出了能力，您应该会收到一个启动错误代码，而不是挂起（如果您不确定，可以使用 deviceQuery SDK 代码检查最大尺寸）。

可能是 cuda-memcheck 正试图获得对 GPU 的独占控制权，并且正在超时，因为其他东西正在使用它 [e.g.你的 X 服务器] - cuda-gdb 工作得更好吗，这些工具是否适用于其他代码？

Answer 2

您的代码运行时间太长（是的，超过 60 秒）。

即使你的内核“什么都不做”，它仍然代表一个__global__ 函数调用。为方便起见，编译器生成了大量的前导代码。通常，编译器会优化大部分前导代码，因为您的函数什么都不做（例如，它对传递给它的变量不做任何事情，前导代码使每个线程都可以使用该变量。）但是，当您传递-G 开关时，您消除了几乎所有编译器优化。通过获取可执行文件和inspecting the code 和cuobjdump -sass ...，您可以了解每个线程块实际运行的代码大小。

其次，使用cuda-memcheck 运行代码通常会增加执行时间。 cuda-memcheck 执行程序调整顺序并降低线程块的执行速率，因此它可以全面分析每个线程块的内存访问模式等。

最终结果是您的空内核调用，部分原因是网格非常大（需要处理超过 1 亿个线程块），执行时间超过 60 秒。如果您想验证这一点，请将您的 TDR 超时时间增加到 5 分钟或 10 分钟，最终您会看到程序正常返回。

在我的例子中，使用-G 和cuda-memcheck，您的程序需要大约 30 秒才能在具有 11 个 SM 的 Quadro5000 GPU 上运行。您的 cc2.1 GPU 可能有大约 2 个 SM，因此运行速度会比我的还要慢。如果我在没有-G 开关的情况下进行编译，运行时间会下降到大约 2 秒。如果我用-G 开关编译，但没有cuda-memcheck 运行，大约需要4 秒。如果我从内核中删除int a 参数（这大大减少了前导代码），我可以使用-G 编译并使用cuda-memcheck 运行，只需要2 秒。

带有-G和int a参数的内核机器码：

            Function : _Z4stami
    .headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
    /*0000*/         MOV R1, c[0x1][0x100];            /* 0x2800440400005de4 */
    /*0008*/         ISUB R1, R1, 0x8;                 /* 0x4800c00020105d03 */
    /*0010*/         S2R R0, SR_LMEMHIOFF;             /* 0x2c000000dc001c04 */
    /*0018*/         ISETP.GE.AND P0, PT, R1, R0, PT;  /* 0x1b0e00000011dc23 */
    /*0020*/     @P0 BRA 0x30;                         /* 0x40000000200001e7 */
    /*0028*/         BPT.TRAP;                         /* 0xd00000000000c007 */
    /*0030*/         IADD R0, R1, RZ;                  /* 0x48000000fc101c03 */
    /*0038*/         MOV R2, R0;                       /* 0x2800000000009de4 */
    /*0040*/         MOV R3, RZ;                       /* 0x28000000fc00dde4 */
    /*0048*/         MOV R2, R2;                       /* 0x2800000008009de4 */
    /*0050*/         MOV R3, R3;                       /* 0x280000000c00dde4 */
    /*0058*/         MOV R4, c[0x0][0x4];              /* 0x2800400010011de4 */
    /*0060*/         MOV R5, RZ;                       /* 0x28000000fc015de4 */
    /*0068*/         IADD R2.CC, R2, R4;               /* 0x4801000010209c03 */
    /*0070*/         IADD.X R3, R3, R5;                /* 0x480000001430dc43 */
    /*0078*/         MOV32I R0, 0x20;                  /* 0x1800000080001de2 */
    /*0080*/         LDC R0, c[0x0][R0];               /* 0x1400000000001c86 */
    /*0088*/         IADD R2.CC, R2, RZ;               /* 0x48010000fc209c03 */
    /*0090*/         IADD.X R3, R3, RZ;                /* 0x48000000fc30dc43 */
    /*0098*/         MOV R2, R2;                       /* 0x2800000008009de4 */
    /*00a0*/         MOV R3, R3;                       /* 0x280000000c00dde4 */
    /*00a8*/         ST.E [R2], R0;                    /* 0x9400000000201c85 */
    /*00b0*/         BRA 0xc8;                         /* 0x4000000040001de7 */
    /*00b8*/         EXIT;                             /* 0x8000000000001de7 */
    /*00c0*/         EXIT;                             /* 0x8000000000001de7 */
    /*00c8*/         EXIT;                             /* 0x8000000000001de7 */
    /*00d0*/         EXIT;                             /* 0x8000000000001de7 */
            .........................

带有-G但没有int a参数的内核机器码：

            Function : _Z4stamv
    .headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
    /*0000*/         MOV R1, c[0x1][0x100];  /* 0x2800440400005de4 */
    /*0008*/         BRA 0x20;               /* 0x4000000040001de7 */
    /*0010*/         EXIT;                   /* 0x8000000000001de7 */
    /*0018*/         EXIT;                   /* 0x8000000000001de7 */
    /*0020*/         EXIT;                   /* 0x8000000000001de7 */
    /*0028*/         EXIT;                   /* 0x8000000000001de7 */
            .........................