Ở bài viết này mình sẽ viết tiếp về cách sử dụng NVIDIA Compute Sanitizer, hãy đọc những bài này: NVIDIA Compute Sanitize phần 1, Data Hazard trước khi đọc bài viết này

NVIDIA Compute Sanitizer

Nối tiếp ở phần 1, phần 2 này sẽ đi qua 2 cái còn lại là:

• Racecheck, a shared memory data access hazard detection tool
• Synccheck for thread synchronization hazard detection

Racecheck

theo như nvidia đã đề cập về NVIDIA Compute Sanitizer thì cái racecheck dùng để check khi sử dụng trên shared memory nên nếu các bạn test trên global memory thì sẽ không có kết quả nha

Code

__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 code này y chang đoạn code ở bài data hazard và đây là nguyên lí hoạt động của nó

chỉ khác 1 chỗ là thay vì chúng ta xử dụng global memory thì ở đây chúng ta dùng shared memory

Shared memory mình sẽ nói riêng vì đây là phần kiến thức rất quan trọng khi nói về cuda nên ở bài viết này bạn chỉ cần hiểu là thay vì dùng global memory để thực hiện phép cộng 1+2+3+4 bằng 1 cách song song thì dùng shared memory

và bây giờ chúng ta dùng NVIDIA Compute Sanitizer để check có bị data hazard hay không bằng dòng lệnh

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

Ở đây bạn sẽ thấy 1 điều bất ngờ là kết quả vẫn ra đúng mặc dù đã bị data hazard, như ở bài trước mình đã nói về hiện tượng "undifined behavior" thì đây chính là nó, chúng ta sẽ không thể xác định được liệu nó có gây ra lỗi hay là không, có thể máy mình ra kết quả đúng NHƯNG máy của các bạn thì lại khác ==> nên hiện tượng "undifined behavior" khá là phiền phức

Lúc này nếu chúng ta dùng __syncthread() thì sẽ giải quyết vấn đề này

Còn về Synccheck thì theo mình test cũng như tìm các tư liệu trên nvidia blog thì thấy họ không đề cập gì về code hết nên mình cũng không thể minh họa cho các bạn nên mình sẽ bỏ qua, nếu mình tìm được thì mình sẽ bổ sung sau ( nếu các bạn tìm thấy thì hãy bình luận ở dưới nha)

Bài tập

Đúng ra lỗi của chúng ta chỉ là 4 data hazard ( N =4 ) nhưng tại sao bức ảnh ở phía trên chúng ta lại bị 2 lỗi data hazard ( 4 và 8 )

Gợi ý: 1 data hazard = 1read or 1 write

4 data hazard là 4 read hay 4 write?

8 data hazard là 8 read hay 8 write hay 4 read và 4 write