超时后在cudaDeviceSynchronize中终止启动

Question

我试着用三维网格运行一个简单的程序，但是由于某种原因，当我用cuda启动它时，它就会被卡住，超时之后它就被终止了。这个问题与短暂的超时无关，因为我只是用这种方式将其更改为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

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

即使您的内核“什么都不做”，它仍然表示一个__global__函数调用。为了方便这一点，编译器生成了相当数量的序言代码。通常，编译器会优化大部分前导代码，因为您的函数什么也不做(例如，它对传递给它的变量什么也不做，序言代码提供给每个线程)。但是，当您通过-G开关时，几乎消除了所有编译器优化。通过将可执行文件和检查代码与cuobjdump -sass ...一起使用，您可以了解每个线程块实际运行的代码的大小。

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

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

在我的例子中，使用-G和cuda-memcheck，您的程序在Quadro5000 GPU上运行大约需要30秒，GPU有11条短信。您的cc2.1GPU可能有大约2短消息，因此将运行甚至比我的慢。如果我不使用-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 */
            .........................

Answer 2

我刚刚在code和cuda下运行了C2050 (Performance2.0)上的代码，没有问题。这表明，与代码本身相比，这个问题更可能与您的卡/设置有关。

如果您超出了能力，您应该得到一个启动错误代码，而不是挂起(如果您不确定，可以使用deviceQuery SDK代码检查最大大小)。

可能是试图获得GPU的独占控制权，并且由于其他东西正在使用它，比如您的X服务器--cuda是否工作得更好，这些工具是否适用于其他代码？