cuda

Start by adding printf in device code guarded by thread index, then run compute-sanitizer --tool memcheck to catch memory errors. If still stuck, use cuda-gdb -batch -ex "run" -ex "bt" for a backtrace, compiled with -g -G -lineinfo.

Run nsys profile -o report ./program first to see the timeline and identify which kernels dominate runtime. Then run ncu --kernel-name "myKernel" --set basic ./program to analyze that kernel's memory throughput, occupancy, and compute utilization.

nsys answers where time is spent across the whole system: CPU/GPU interaction, kernel launches, and memory transfers. ncu answers why a specific kernel is slow, providing detailed metrics like occupancy, coalescing, bank conflicts, and roofline analysis.

An "Invalid __shared__ write out of bounds" error usually means insufficient dynamic shared memory was allocated in the kernel launch configuration, not wrong array indexing. Check the third parameter in the <<<grid, block, smem_size>>> launch syntax.

Profile with ncu using the l1tex__data_bank_conflicts_pipe_lsu_mem_shared metrics divided by wavefront counts; more than one conflict per operation indicates a problem. Fixes include padding shared memory arrays or changing thread-to-data mapping.

Use printf first, before any tool, since it works when debuggers produce inscrutable output. Guard output with conditions like if (idx == 0) to avoid flooding, and print at kernel entry and suspected failure points.