Giới Thiệu Về NVIDIA Compute Sanitizer

Trong bài viết này, mình sẽ hướng dẫn các bạn cách sử dụng NVIDIA Compute Sanitizer, một công cụ hữu ích giúp các lập trình viên mới bắt đầu với CUDA dễ dàng phát hiện lỗi trong mã nguồn của họ.

Mặc dù đối với những bạn đã quen thuộc với CUDA thì tính năng của NVIDIA Compute Sanitizer có thể không quá cần thiết, nhưng việc nắm bắt kiến thức về công cụ này vẫn rất quan trọng.

Tính Năng Của NVIDIA Compute Sanitizer

NVIDIA Compute Sanitizer hỗ trợ kiểm tra bốn loại lỗi phổ biến mà các lập trình viên CUDA mới thường mắc phải:

• Memcheck: Kiểm tra lỗi truy cập bộ nhớ và phát hiện rò rỉ bộ nhớ.
• Racecheck: Công cụ phát hiện lỗi truy cập dữ liệu trong bộ nhớ chia sẻ.
• Initcheck: Phát hiện lỗi truy cập bộ nhớ toàn cục chưa được khởi tạo.
• Synccheck: Kiểm tra lỗi đồng bộ hóa giữa các luồng.

Ví Dụ Code Cơ Bản

Dưới đây là một đoạn mã đơn giản thực hiện cộng hai vector để minh họa bốn trường hợp lỗi:

#include <stdio.h>

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    c[tid] = a[tid] + b[tid];
}

int main() {
    int n = 10;
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = n * sizeof(int);

    a = (int*)malloc(size);
    b = (int*)malloc(size);
    c = (int*)malloc(size);

    for (int i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i;
    }

    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);

    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

    vectorAdd<<<1, ?>>>(d_a, d_b, d_c, n);

    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

    free(a);
    free(b);
    free(c);

    return 0;
}

Khi bắt đầu lập trình CUDA, nhiều bạn có thể mắc lỗi cơ bản là sử dụng số luồng ít hoặc nhiều hơn dữ liệu, dẫn đến hiện tượng “undefined behavior” mà không có thông báo lỗi cụ thể. Điều này có thể gây ra lỗi logic trong chương trình lớn hơn và ảnh hưởng nghiêm trọng đến việc phân bổ bộ nhớ.

Phát Hiện Lỗi Với Initcheck

Giả sử chúng ta chỉnh sửa mã lệnh như sau để sử dụng 9 luồng:

vectorAdd<<<1, 9>>>(d_a, d_b, d_c, n);

Với N = 10 và 9 luồng, chúng ta sẽ gặp phải vấn đề rò rỉ bộ nhớ. Để kiểm tra, chúng ta sử dụng câu lệnh sau:

compute-sanitizer --tool initcheck --track-unused-memory yes --show-backtrace no

Vì chúng ta sử dụng 10 phần tử (4 bytes mỗi phần tử) nhưng chỉ dùng 9 luồng, khoản bộ nhớ không được sử dụng tương ứng 10% tổng bộ nhớ.

Khi sửa mã lại thành dùng 10 luồng, chúng ta sẽ khắc phục được lỗi này.

Kiểm Tra Lỗi Với Memcheck

Bây giờ chuyển sang trường hợp sử dụng nhiều hơn 10 luồng:

vectorAdd<<<1, 11>>>(d_a, d_b, d_c, n);

Với 11 luồng, ta chỉ khởi tạo đủ cho 10 luồng, dẫn đến lỗi threads out-of-bounds array access. Lệnh kiểm tra lỗi sẽ là:

compute-sanitizer --tool memcheck --show-backtrace no

Khi sao chép qua GPU bằng cudaMemcpy, phần tử thứ 11 sẽ bị lỗi vì dữ liệu không đủ. Để khắc phục, chúng ta cần giới hạn luồng như sau:

if (tid < n) {
    c[tid] = a[tid] + b[tid];
}

Ngoài ra, đừng quên sử dụng cudaFree để tránh lỗi rò rỉ bộ nhớ, lệnh kiểm tra sẽ là:

compute-sanitizer --tool memcheck --leak-check=full --show-backtrace no

Kết Luận

Hai lỗi còn lại là Synccheck và Racecheck sẽ được trình bày chi tiết trong phần sau của bài viết. Hãy theo dõi để không bỏ lỡ kiến thức quan trọng này!
source: viblo