Hướng Dẫn Sử Dụng NVIDIA Compute Sanitizer - Phần 2

Trong bài viết này, chúng ta sẽ tiếp tục khám phá cách sử dụng NVIDIA Compute Sanitizer, đặc biệt là hai công cụ quan trọng: Racecheck và Synccheck. Nếu bạn chưa đọc, hãy tham khảo bài viết NVIDIA Compute Sanitizer Phần 1 và bài viết về Data Hazard để có cái nhìn tổng quan hơn.

NVIDIA Compute Sanitizer

Phần 2 này sẽ tập trung vào hai công cụ sau:

Racecheck: Công cụ phát hiện lỗi truy cập dữ liệu trong bộ nhớ chia sẻ.
Synccheck: Công cụ phát hiện lỗi đồng bộ hóa giữa các luồng.

Công Cụ Racecheck

Theo thông tin từ NVIDIA, Racecheck được sử dụng để kiểm tra các lỗi truy cập dữ liệu khi làm việc với bộ nhớ chia sẻ (shared memory). Vì vậy, nếu bạn thực hiện kiểm tra trên bộ nhớ toàn cục (global memory), bạn sẽ không nhận được kết quả.

Mã Nguồn

cpp Copy

__global__ void sumWithSharedMemory(int* input) {
    __shared__ int sharedMemory[4];

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    sharedMemory[tid] = input[i];

    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        // __syncthreads(); -----> barrier

        if (tid % (2 * stride) == 0) {
            sharedMemory[tid] += sharedMemory[tid + stride];
        }
    }

    printf("blockIdx.x=%d --> %d\n", blockIdx.x, sharedMemory[tid]);
}

Đoạn mã này tương tự như đoạn mã trong bài viết về Data Hazard, nhưng điểm khác biệt lớn là ở đây chúng ta sử dụng bộ nhớ chia sẻ thay vì bộ nhớ toàn cục.

Ghi chú về bộ nhớ chia sẻ: Đây là kiến thức cực kỳ quan trọng trong CUDA, nên trong bài viết này, bạn chỉ cần hiểu rằng bộ nhớ chia sẻ giúp thực hiện phép cộng song song hiệu quả hơn.

Sử Dụng NVIDIA Compute Sanitizer

Để kiểm tra các lỗi dữ liệu, bạn có thể sử dụng dòng lệnh sau:

Copy

compute-sanitizer --tool racecheck --show-backtrace no ./a.out

Điều thú vị là kết quả vẫn cho ra đúng mặc dù có sự cố dữ liệu. Như đã đề cập trước đó về hiện tượng "undefined behavior", điều này có nghĩa là chúng ta không thể chắc chắn liệu sự cố có gây ra lỗi hay không. Kết quả có thể đúng trên máy của bạn nhưng lại khác trên máy khác, vì vậy hiện tượng "undefined behavior" thật sự rất phiền phức.

Nếu chúng ta sử dụng __syncthreads() trong mã, điều này sẽ giải quyết được vấn đề này.

Công Cụ Synccheck

Tôi đã thử nghiệm cũng như tìm kiếm tài liệu trên blog NVIDIA nhưng không tìm thấy mã ví dụ nào cho Synccheck. Do đó, mình sẽ bỏ qua phần này và nếu tìm được thông tin cụ thể sau này, mình sẽ cập nhật. Nếu bạn có thông tin nào, hãy bình luận bên dưới nhé!

Bài Tập

Trong thực tế, lỗi của chúng ta chỉ là 4 sự cố dữ liệu (N = 4). Tuy nhiên, tại sao kết quả lại hiện ra 2 lỗi dữ liệu (4 và 8)?

Gợi ý: 1 sự cố dữ liệu = 1 lần đọc hoặc 1 lần ghi.
4 sự cố dữ liệu có phải là 4 lần đọc hay 4 lần ghi?
8 sự cố dữ liệu có phải là 8 lần đọc hay 8 lần ghi, hay 4 lần đọc và 4 lần ghi?

source: viblo

[NVIDIA Tools] Hướng Dẫn Sử Dụng NVIDIA Compute Sanitizer - Phần 2: Phát Hiện Lỗi Dữ Liệu và Đồng Bộ Hóa