Ở bài viết này mình sẽ hướng dẫn các bạn sử dụng NVIDIA Compute Sanitizer, 1 công cụ rất tuyệt vời để hỗ trợ cho các bạn mới bắt đầu về cuda

Đối với các bạn đã quá quen thuộc về cuda thì NVIDIA Compute Sanitizer không giúp ích được mấy nhưng dù sao biết vẫn hơn không

NVIDIA Compute Sanitizer

NVIDIA Compute Sanitizer giúp chúng ta check 4 lỗi quan trọng mà các cuda beginner hay mắc phải:

• Memcheck for memory access error and leak detection
• Racecheck, a shared memory data access hazard detection tool
• Initcheck, an uninitialized device global memory access detection tool
• Synccheck for thread synchronization hazard detection

Đây là 1 đoạn code đơn giản ( cộng 2 vector) để phân tích 4 trường hợp

#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 mới bắt đầu code cuda, các bạn sẽ mắc vào 1 lỗi cơ bản là sử dụng thread ít/nhiều hơn data và điều này sẽ dẫn tới 1 bug “undefined behavior”, Nó có thể hoạt động mà không có báo lỗi. Trong một chương trình lớn hơn, nó có thể gây ra các vấn đề về logic và nghiêm trọng ảnh hưởng đến việc phân bổ các memory khác

Initcheck

Ở đây mình sẽ ví dụ cụ thể

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

Như các bạn có thể thấy, chúng ta có N = 10 nhưng chỉ dùng 9 thread để xử lí và việc này dẫn đến leak memory, và chúng ta sẽ dùng câu lệnh này để profile

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

Chúng ta dùng N =10 (int) nên tổng byte là 40 bytes và sử dụng 9 thread ==> còn dư 10% memory không được sử dụng

Và đây là kết quả khi ta sửa lại thành dùng 10 thread

Memcheck

Ở phía trên là trường hợp sử dụng thread ít hơn thì bây giờ chúng ta đến trường hợp thread nhiều hơn

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

Như các bạn có thể thấy, chúng ta chỉ khởi tạo đủ cho 10 thread hoạt động ==> dẫn tới thread thứ 11 sẽ bị threads out-of-bounds array access và chúng ta sẽ dùng câu lệnh này để profile

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

Để giải thích đơn giản là khi chúng ta dùng cudaMemcpy để copy qua GPU thì phần tử thứ 11 sẽ bị fail vì khi khởi tạo chỉ đủ cho 10 phần tử ==> và việc fail này là thread(10,0,0 ) tức là thread thứ 11 sẽ bị nằm ngoài rìa dẫn đến “undefined behavior”

Các khắc phục là chúng ta chỉnh về dùng 10 thread hoặc thêm boundary

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

Và đây là kết quả

Ngoài ra nếu các bạn để ý thì đoạn code này chúng ta thiếu cudaFree nên cũng sẽ dẫn tới leak memory và chúng ta sẽ dùng câu lệnh này để profile

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

2 lỗi còn lại ( Synccheck - Racecheck) mình sẽ nói sau khi hết bài atomic function và data hazard