Lập trình song song: Bài 11 - Data Hazard

Khi bàn về lập trình song song, một trong những hiện tượng quan trọng mà chúng ta cần lưu ý là data hazard. Đây là một loại lỗi logic mà các lập trình viên thường gặp phải trong quá trình phát triển ứng dụng song song. Tuy nhiên, với sự hỗ trợ của công cụ NVIDIA Compute Sanitizer, việc phát hiện và khắc phục những lỗi này đã trở nên dễ dàng hơn. Trong bài viết này, chúng ta sẽ cùng tìm hiểu về data hazard, cách nó xảy ra và làm thế nào để khắc phục.

Lưu ý: Để hiểu rõ hơn về vấn đề này, mình khuyên bạn nên đọc bài viết Synchronization - Asynchronization trước khi tiếp tục.

Data Hazard Là Gì?

Data hazard là hiện tượng xảy ra khi nhiều thread đồng thời đọc và ghi đến cùng một giá trị, dẫn đến tình trạng xung đột dữ liệu. Có hai vấn đề chính liên quan đến data hazard:

1. Data Race: Đây là hiện tượng xảy ra khi hai hoặc nhiều thread đồng thời truy cập một biến mà không có sự đồng bộ hóa. Cụ thể, nó thường liên quan đến các thao tác write after read (ghi sau khi đọc) hoặc read after write (đọc sau khi ghi). Khi không có đồng bộ hóa, một thread có thể ghi đè lên dữ liệu mà thread khác đang sử dụng, dẫn tới xung đột giá trị.

2. Race Condition: Khái niệm này sâu hơn và không chỉ giới hạn ở việc 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 bị ảnh hưởng bởi các sự kiện không xác định, thường được gọi là undefined behavior.

Như vậy, khi lập trình CUDA, bạn luôn phải chú ý đến việc nhiều thread có thể đồng thời truy cập vào cùng một giá trị để đảm bảo tính chính xác trong chương trình.

Minh Họa Về Data Hazard

Dưới đây là một ví dụ code sử dụng CUDA để minh họa hiện tượng data hazard:

#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;
}

Trong đoạn mã trên, chúng ta thấy rằng không có sự đồng bộ giữa các thread, dẫn đến hiện tượng data race. Cụ thể, tại bước 1, giá trị 3+4 chưa được hoàn thành thì đã chuyển sang bước 2, dẫn tới kết quả sai (3 + 3 = 6 thay vì 3 + 7 = 10).

Giải Quyết Vấn Đề Data Hazard

Để khắc phục vấn đề này, chúng ta cần sử dụng câu lệnh __syncthreads() như một barrier để các thread có thể chờ nhau hoàn tất trước khi tiếp tục thực hiện các thao tác tiếp theo. Như vậy, việc đồng bộ hóa sẽ giúp tránh được hiện tượng xung đột dữ liệu và đảm bảo kết quả chính xác hơn trong quá trình thực hiện.

Qua bài viết này, hi vọng bạn đã có cái nhìn rõ hơn về data hazard trong lập trình song song và cách để phòng tránh nó.
source: viblo