Untitled

# cat t228.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 t228 t228.cu -g -G -arch=sm_89
ptxas warning : Stack size for entry function '_Z6kernelv' cannot be statically determined
# cuda-gdb ./t228
NVIDIA (R) CUDA Debugger
CUDA Toolkit 12.2 release
Portions Copyright (C) 2007-2023 NVIDIA Corporation
GNU gdb (GDB) 12.1
Copyright (C) 2022 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 GDB was configured as "x86_64-pc-linux-gnu".
Type "show configuration" for configuration details.
For bug reporting instructions, please see:
<https://www.gnu.org/software/gdb/bugs/>.
Find the GDB manual and other documentation resources online at:
    <http://www.gnu.org/software/gdb/documentation/>.

For help, type "help".
Type "apropos word" to search for commands related to "word"...
Reading symbols from ./t228...
(cuda-gdb) b 11
Breakpoint 1 at 0xabf8: file /root/bobc/t228.cu, line 20.
(cuda-gdb) b 14
Note: breakpoint 1 also set at pc 0xabf8.
Breakpoint 2 at 0xabf8: file /root/bobc/t228.cu, line 20.
(cuda-gdb) r
Starting program: /root/bobc/t228
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[New Thread 0x7ffff55a9000 (LWP 426949)]
[Detaching after fork from child process 426950]
[New Thread 0x7ffff4a2c000 (LWP 426968)]
[New Thread 0x7fffe9fff000 (LWP 426969)]
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]

Thread 1 "t228" hit Breakpoint 1, device_func (octants=0x7fffed000000, level=0) at t228.cu:11
11              octants[0] = 0; //! breakpoint 1
(cuda-gdb)