Khi chúng ta nhắc đến song song chúng ta sẽ nhắc tới hiện tượng data hazard, 1 bug khiến chúng ta khá là đau đầu khi fix vì đây là lỗi về mặt logic NHƯNG bây giờ chúng ta đã có công cụ NVIDIA Compute Sanitizer nên việc fix bug này cũng trở nên đỡ hơn 1 phần nào. Bài viết này mình sẽ nói về data hazard là gì và minh họa nó

Sẽ tuyệt hơn nếu các bạn đọc bài Synchronization - Asynchronization trước khi đọc bài này

Data Hazard

Hiện tượng các thread read và write 1 value nào đó thì sẽ dẫn đến xung đột và hiện tượng đó gọi là data hazard

Khi nói về data hazard chúng ta sẽ đụng đến 2 vấn đề:

• Data Race: Thường liên quan đến việc "write after read" hoặc "read after write", nhưng nó chủ yếu tập trung vào việc đồng thời truy cập (đọc và/hoặc ghi) một biến lưu trữ mà không có sự đồng bộ hóa. Điều này có thể dẫn đến việc 1 thread ghi đè lên dữ liệu mà thread khác đang đọc hoặc chuẩn bị ghi ==> dẫn đến conflict data value
• Race condition: Khái niệm này rộng hơn và không chỉ giới hạn ở truy cập dữ liệu. Race condition xảy ra khi kết quả cuối cùng của một hệ thống xuất hiện sự kiện không xác định hay còn gọi là "undefined behavior"

Tóm lại các bạn chỉ cần nhớ: khi code cuda cần lưu ý hiện tượng các thread cùng truy cập vào cùng 1 value để xử lí

Minh họa

#include <stdio.h>
#include <cuda_runtime.h>

#define ARRAY_SIZE 4


__global__ void sum(int *d_array)
{
    int id = blockIdx.x * blockDim.x + threadIdx.x;

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

        if (threadIdx.x % (2 * stride) == 0)
        {
            d_array[id] += d_array[id + stride];
        }
    }
    printf("blockIdx.x=%d --> %d\n", blockIdx.x, d_array[id]);
}

int main()
{
    int h_array[4] = {1, 2, 3, 4};
    int *d_array;

   cudaMalloc((void **)&d_array, sizeof(int) * ARRAY_SIZE);
   cudaMemcpy(d_array, h_array, sizeof(int) * ARRAY_SIZE, cudaMemcpyHostToDevice);

    sum<<<1, 4>>>(d_array);

    cudaFree(d_array);

    return 0;
}

Đây là hình ảnh nguyên lí code hoạt động

Ở đây chúng ta không đồng bộ các thread dẫn tới data race ( vào step 1: khi 3+4 chưa xong thì đã qua step 2 nên 3 + 3 =6 chứ không phải 3+ 7 =10 )

Để giải quyết vấn đề này thì chúng ta chỉ cần đặt 1 barrier để cho các thread đợi lẫn nhau cho đến khi các thread chậm nhất xong bằng câu lệnh __syncthreads()

và output khi thêm syncthreads