CUDA Tools: cuda-gdb

这是CUDA以及GPU学习的第二个系列（上一个系列是GPU后端优化）。在CUDA学习中，高效的调试手段以及性能分析手段能够辅助我们编写出正确高性能的cuda函数。作为本系列的第一章，我们先来了解一下cuda如何debug，学习cuda-gdb工具的使用。

最简CUDA 项目框架

这里提供一个用cmake构建cuda项目的最简框架，以归约算子的cuda实现为例。我们的文件目录如下：

├── CMakeLists.txt
├── cuda_reduce_study
│   ├── CMakeLists.txt
│   ├── HelloWorld.cu
├── README.md
└── run.sh

顶层的CMakeLists如下：

cmake_minimum_required(VERSION 3.30)
project(cuda_practice VERSION 0.1.0 LANGUAGES CXX C CUDA)

# 设置可执行文件输出目录（构建目录下的bin文件夹）
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin)

find_package(CUDAToolkit)
add_subdirectory(cuda_reduce_study)

下一层CMakeLists如下：

cmake_minimum_required(VERSION 3.18)
project(cuda_reduce_study LANGUAGES CUDA CXX)

# Debug 模式下添加 CUDA 调试选项
if (CMAKE_BUILD_TYPE STREQUAL "Debug")
    message(STATUS "Enabling CUDA debug flags: -G -O0 -g")
    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -G -O0 -g")
endif()

# ============ Hello World ============ #
add_executable(hello_world HelloWorld.cu)
target_link_libraries(hello_world PRIVATE CUDA::cudart CUDA::cublas)

这里需要注意的是，在CUDA编译器nvcc中，需要-G -oO -g才能成功开启debug标志。

通过如下构建脚本，即可构建运行：

#!/bin/bash

if [ -d "build" ]; then
	rm -rf build
fi

mkdir build
cd build
cmake .. \
	-DCMAKE_BUILD_TYPE=Debug \
	-DCMAKE_EXPORT_COMPILE_COMMANDS=1
make -j $(nproc)

VSCode

插件配置

debug需要在extension处配置Nsight Visual Studio Edition。

Launch.json

需要配置如下launch.json：

{
    // Use IntelliSense to learn about possible attributes.
    // Hover to view descriptions of existing attributes.
    // For more information, visit: https://go.microsoft.com/fwlink/?linkid=830387
    "version": "0.2.0",
    "configurations": [
        {
            "name": "CUDA C++: launch",
            "type": "cuda-gdb",
            "request": "launch",
            "program": "",
        },
        {
           "name": "CUDA C++: launch",
            "type": "cuda-gdb",
            "request": "attach", 
        }
    ]
}

CUDA-GDB使用指南

一种方法是运行cuda-gdb命令行，另一种是在vscode通过F5 debug。以reductionShuffle为例，测试代码如下：

#include <bits/stdc++.h>
#include <cstdlib>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <sys/time.h>
#include "device_launch_parameters.h"
#include "cuda_runtime.h"

# define THREAD_PER_BLOCK 256
#define FULL_MASK 0xffffffff

// 核函数：规约函数，支持shuffle的版本
template <unsigned int BLOCK_SIZE, int TASK_PER_THREAD>
__global__ void reduce(float *device_input, float *device_output) {
    // 计算每个快的索引位置
    float *device_data = device_input + blockIdx.x * TASK_PER_THREAD * BLOCK_SIZE;     // 每个block的起始位置，计算逻辑是一个thread可以处理多少任务，一个block有多少个thread
    // 由此计算出一个block可以计算多少任务，每次位移这么多

    // 用register
    float sum = 0.0;
    // 这里需要循环，做multiAdd
    for (int  i = 0; i < TASK_PER_THREAD; i++) {
        sum += device_data[i * BLOCK_SIZE + threadIdx.x]; // 每个线程计算自己的数据
    }
   
    // 一个warp的计算，利用shuffle instruction
    // 具体参考https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/
    for (int offset = 16; offset > 0; offset /= 2) {
        sum += __shfl_down_sync(FULL_MASK, sum, offset); // 规约
    } 

    // 在一个warp内归约好后，还要汇总各个warp的结果。由于GPU最多1024个线程，所以这里的warp数不会超过32，所以最多两层warp内reduce
    // 这里需要用shared memory来缓存结果
    __shared__ float shared_data[32];
    unsigned int lane_id = threadIdx.x % warpSize; // 计算lane id
    unsigned int warp_id = threadIdx.x / warpSize; // 计算warp id
    // 只有每个warp的第一个lane参与计算
    if (lane_id == 0) {
        shared_data[warp_id] = sum; // 将结果写入共享内存
    }
    __syncthreads(); // 同步线程，确保所有线程(按照warp执行)都完成数据写入

    // 只保留第一个线程的结果
    if (threadIdx.x == 0) {
        device_output[blockIdx.x] = sum; // 将结果写入输出
    }

    // 继续第二轮的reduction
    if (warp_id == 0) {
        // 从共享内存读入有个trick，就是可能32个缓存中，后面的warp没有写入数据，均是零
        // 加入有128个线程per block，则warp有4个，所以到了第二轮，只有前4个需要计算，后面的共享内存全为零
        sum = (lane_id < (blockDim.x / 32)) ? shared_data[lane_id] : 0.0; // 读取共享内存中的结果\

        for (int offset = 16; offset > 0; offset /= 2) {
            sum += __shfl_down_sync(FULL_MASK, sum, offset); // 规约
        } 
    }

    if (threadIdx.x == 0) {
        device_output[blockIdx.x] = sum; // 将结果写入输出
    }
}

// 检查结果正确性
bool check(float *out, float *res, int n) {
    for (int i = 0; i < n; i++) {
        if (fabs(out[i] - res[i]) > 0.05) {       // 我们允许的误差范围
            printf("Error: out[%d] = %f, res[%d] = %f\n", i, out[i], i, res[i]);
            return false;
        }
    }
    return true;
}

int main() {
    const int N = 32 * 1024 * 1024; // 总计算量

    //========== 输入内存分配 ==============
    float *host_input = (float *)malloc(N * sizeof(float));       // 主机输入
    float *device_input; // 设备输入
    cudaMalloc((void **)&device_input, N * sizeof(float)); // 设备输入，注意传入二级指针

    //========== 输出内存分配 ==============
    constexpr int block_num = 1024; // 指定block num是1024
    constexpr int task_per_block = N / block_num; // 每个block的任务数
    constexpr int task_per_thread = task_per_block / THREAD_PER_BLOCK; // 每个线程的任务数

    float *host_output = (float *)malloc(block_num * sizeof(float));       // 主机输出
    float *device_output; // 设备输出
    cudaMalloc((void **)&device_output, block_num * sizeof(float)); // 设备输出，注意传入二级指针
    float *res = (float *)malloc(block_num * sizeof(float));                // 结果

    //========== 输入数据初始化 ==============
    for (int i = 0; i < N; i++) {
        // 生成随机数l
        host_input[i] = 2.0 * (float)drand48() - 1.0; // [-1, 1]之间的随机数
    }

    //========== CPU计算 ==============
    for (int i = 0; i < block_num; i++) {
        res[i] = 0.0;
        for (int j = 0; j < task_per_block; j++) {
            res[i] += host_input[i * task_per_block + j];
        }
    }

    //=========== GPU计算 ==============
    // 1. 将数据从主机拷贝到设备
    cudaMemcpy(device_input, host_input, N * sizeof(float), cudaMemcpyHostToDevice);

    // 2. 启动核函数
    dim3 Grid(block_num, 1);       
    dim3 Block(THREAD_PER_BLOCK, 1);
    reduce<THREAD_PER_BLOCK, task_per_thread><<<Grid, Block>>>(device_input, device_output);

    // 3. 将结果从设备拷贝到主机
    // 目前的reduce是每个block计算一个值，所以我们需要将每个block的结果拷贝到主机
    cudaMemcpy(host_output, device_output, block_num * sizeof(float), cudaMemcpyDeviceToHost);

    // 检测结果
    if (check(res, host_output, block_num)) {
        printf("================== Reduce Shuffle ==================\n");
        printf("Result is correct!\n");
    } else {
        printf("================== Reduce Shuffle ==================\n");
        printf("Result is incorrect!\n");
    }

    cudaFree(device_input);
    cudaFree(device_output);
    free(host_input);
    free(host_output);
    free(res);

    return 0;
}

相比常用的gdb，cuda-gdb最大的特点是能够选择跟踪的block和thread：

有这个功能的支持，我们能够跟踪各个线程的逻辑。

参考资料

1. CUDA-GDB User Manual