CUDA 递归的奇怪行为

Question

在下面的最小可重现示例中，当

device_func

中的递归处于活动状态时，

__synchthreads()

屏障将被忽略，并且在调试时，断点 2 发生在断点 1 之前。如果删除递归，则它将按预期工作。怎么会这样？该代码是使用 CUDA Toolkit 12.5 为 NVIDIA Quadro P600 使用

nvcc -arch=sm_61 -G -g example.cu

进行编译的。

#include <cstdint>
#include <cstddef>
#include <cstdio>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

__device__ void device_func(uint32_t *octants, size_t const level)
{
    if (threadIdx.x == 0)
    {
        octants[0] = 0; //! breakpoint 1
    }
    __syncthreads();
    printf("first octant: %d\n", octants[0]); //! breakpoint 2
    //! if used, breakpoint 2 occurs before breakpoint 1
    if (level < 2)
    {
        device_func(octants, level + 1);
    };
}

__global__ void kernel()
{
    __shared__ uint32_t octants[9];

    if (threadIdx.x == 0)
    {
        octants[0] = 99999;
    }

    __syncthreads();
    device_func(octants, 0);
    return;
}

void checkCudaError(cudaError_t err, const char *msg)
{
    if (err != cudaSuccess)
    {
        fprintf(stderr, "CUDA error at %s: %s\n", msg, cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
}

int main()
{
    kernel<<<1, 32>>>();

    // Check for kernel launch errors
    checkCudaError(cudaGetLastError(), "Kernel launch");

    // Wait for the kernel to finish executing
    checkCudaError(cudaDeviceSynchronize(), "Kernel execution");

    return 0;
}

这是调试会话：

$ nvcc -o example -arch=sm_61 -G -g example.cu 
ptxas warning : Stack size for entry function '_Z6kernelv' cannot be statically determined
$ cuda-gdb ./example
NVIDIA (R) cuda-gdb 12.5
Portions Copyright (C) 2007-2024 NVIDIA Corporation
Based on GNU gdb 13.2
Copyright (C) 2023 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.
Type "show copying" and "show warranty" for details.
This CUDA-GDB was configured as "x86_64-pc-linux-gnu".
Type "show configuration" for configuration details.
For bug reporting instructions, please see:
<https://forums.developer.nvidia.com/c/developer-tools/cuda-developer-tools/cuda-gdb>.
Find the CUDA-GDB manual and other documentation resources online at:
    <https://docs.nvidia.com/cuda/cuda-gdb/index.html>.

For help, type "help".
Type "apropos word" to search for commands related to "word"...
Reading symbols from ./example...
(cuda-gdb) b 11
Breakpoint 1 at 0xabfc: file /home/marco/phase_2_barnes_hut/example.cu, line 20.
(cuda-gdb) b 14
Note: breakpoint 1 also set at pc 0xabfc.
Breakpoint 2 at 0xabfc: file /home/marco/phase_2_barnes_hut/example.cu, line 20.
(cuda-gdb) r
Starting program: /home/marco/phase_2_barnes_hut/example 
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[New Thread 0x7ffff6dff000 (LWP 1519)]
[New Thread 0x7ffff5b21000 (LWP 1520)]
[Detaching after fork from child process 1521]
[New Thread 0x7ffff5320000 (LWP 1529)]
[Thread 0x7ffff5320000 (LWP 1529) exited]
[New Thread 0x7ffff5320000 (LWP 1530)]
[New Thread 0x7fffe1fff000 (LWP 1532)]
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (1,0,0), device 0, sm 0, warp 0, lane 1]

CUDA thread hit Breakpoint 2, device_func (octants=0x7fffe5000000, level=0) at example.cu:14
14          printf("first octant: %d\n", octants[0]); //! breakpoint 2
(cuda-gdb)

Answer 1

提交错误后，我被告知这是 Pascal GPU 所期望的，因为不支持发散线程。更多信息这里。