Hiểu Biết Về Hàm Atomic Trong CUDA

Trong bài viết này, chúng ta sẽ khám phá một trong những hàm có sẵn rất hữu ích trong CUDA - hàm atomic. Trước khi tiếp cận nội dung chính, mình khuyên bạn nên tìm hiểu qua các bài viết về Data Hazard và Synchronization - Asynchronization để có nền tảng vững chắc hơn.

Hàm Atomic Là Gì?

Thư viện hàm atomic được NVIDIA phát triển nhằm giải quyết vấn đề data hazard trong lập trình song song. Mục đích chính của nó là đồng bộ hóa các thread trong quá trình thực hiện các phép toán, đảm bảo rằng dữ liệu được cập nhật một cách chính xác mà không gặp phải lỗi không mong muốn.

Triển Khai Thí Dụ Đơn Giản

Giả sử chúng ta có đoạn mã sau:

for(int i = 0; i < n; i++) {
    y += x[i];
}

Khi chuyển đoạn mã này sang lập trình song song với CUDA, bạn sẽ có:

int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < n) {
    *d_result += d_data[tid];
}

Tuy nhiên, đoạn mã trên có thể dẫn đến data hazard do không có sự đồng bộ giữa các thread, gây ra kết quả không chính xác. Đó chính là lý do hàm atomic ra đời.

Cơ Chế Hoạt Động Của Hàm Atomic

Hàm atomic hoạt động theo cách tương tự như vòng lặp for: nó sẽ duyệt qua từng thread một
(thread nào khởi động trước sẽ được xử lý trước) KHI VÀ CHỈ KHI nhiều thread cùng thực hiện read-write trên cùng một giá trị. Ở các trường hợp khác, các thread vẫn hoạt động song song như bình thường.

Cách Sử Dụng Hàm Atomic

Dưới đây là cách viết sử dụng hàm atomic trong mã nguồn:

int tid = threadIdx.x + blockDim.x * blockIdx.x;
atomicAdd(result, array_add[tid]);

Các hàm atomic cơ bản mà bạn nên biết bao gồm:

• atomicAdd(result, array_add[tid]);
• atomicSub(result, array_sub[tid]);
• atomicMax(result, array_max[tid]);
• atomicMin(result, array_min[tid]);

Ngoài ra, còn một số hàm atomic thú vị khác như atomicCAS và atomicExch mà bạn có thể tìm hiểu thêm.

Mình sẽ cung cấp toàn bộ mã nguồn để bạn có thể tham khảo và thực hành. Hy vọng bài viết này sẽ giúp bạn hiểu rõ hơn về hàm atomic và ứng dụng của nó trong lập trình song song với CUDA.
source: viblo