Với việc GPU ngày càng được tận dụng trong nhiều tác vụ tính toán khác nhau, đặc biệt là trong AI và Deep Learning, việc biết 1 chút về lập trình với nền tảng CUDA đã trở thành 1 kĩ năng nên có của 1 AI Engineer. Tuy nhiên, sự khác biệt giữa cách xử lý lệnh của GPU và CPU, bugs nghiễm nhiên sẽ luôn xuất hiện trong các đoạn code kernel của ta và không phải ai cũng có khả năng dành lượng lớn thời gian để ngồi mò và debug trên terminal. Chính vì điều đó mà NVIDIA đã tích hợp NVIDIA-toolkit vào các IDE và text editor phổ thông ngày nay như Eclipse, Visual Studio và Visual Studio Code. Chủ đề bài viết của hôm nay sẽ đi vào việc sử dụng extension Nsight Visual Studio Code Edition để debug kernel code ngay trên trình debugger của VSC. Lưu ý: Hiện tại chỉ hỗ trợ debug GPU code trên target system là Linux

Cài đặt môi trường

Trước hết chúng ta cần cài đặt NVIDIA Toolkit. NVIDIA Toolkit bao gồm thư viện giao tiếp với GPU, công cụ debug và optimize, trình biên dịch C/C++ và thư viện runtime.

1. Cài đặt NVIDIA Toolkit theo hướng dẫn: https://developer.nvidia.com/cuda-downloads
2. Cài NVIDIA Nsight Visual Studio Code Edition trên Visual Studio Code Lưu ý: Nếu target system là WSL2 trên Windows, cần kiểm tra và tạo registry key sau: HKEY_LOCAL_MACHINE\SOFTWARE\NVIDIA Corporation\GPUDebugger\EnableInterface sau đó set thành (DWORD) = 1. Nếu chỉ cài NVIDIA Toolkit trên WSL2 distro mà không cài trên Windows thì thiếu mất key trên.

Vector addition

Tiếp theo chúng ta cần đoạn code GPU để compile và debug. Trong khuôn khổ bài viết, ta sẽ dùng đoạn code sau:

__global__ void add(int n, float *x, float *y)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
        y[i] = x[i] + y[i];
}

int main()
{
    int N = 1 << 20;
    float *x, *y;

    cudaMallocManaged(&x, N * sizeof(float));
    cudaMallocManaged(&y, N * sizeof(float));

    for (int i = 0; i < N; i++)
    {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    add<<<numBlocks, blockSize>>>(N, x, y);

    cudaDeviceSynchronize();

    cudaFree(x);
    cudaFree(y);

    return 0;
}

Đoạn code trên thực hiện phép tính cộng trên vector đơn giản. Mỗi phép tính được đưa vào xử lý trên 1 thread trên GPU và kết quả được lưu thẳng vào vector đầu vào thứ 2. Thực hiện compile đoạn code với nvcc, đặt flag debug -g và flag GPU -G để compiler sinh symbol trên cả CPU lẫn GPU code:

nvcc -g -G main.cu -o main

Cài đặt lauch.json trên vscode và bắt đầu debug

Tạo thư mục .vscode và tạo file launch.json với setting sau:

{
    "version": "0.2.0",
    "configurations": [
        {
            "name": "CUDA C++: Launch",
            "type": "cuda-gdb",
            "request": "launch",
            "program": "${workspaceFolder}/main",
        }
    ]
}

Đặt đường dẫn tới file executable compiled bởi nvcc và Ctrl + F5 để bắt đầu debug session thôi. Đến đây là ta có thể bắt đầu đặt breakpoint trong kernel code, recompile và debug như cách ta debug chương trình C/C++ thông thường rồi.