Ở bài này mình sẽ giới thiệu các bạn một cái built in function khá là xịn trong cuda, và 1 lưu ý là hãy đọc các bài này Data Hazard và Synchronization - Asynchronization trước khi đọc bài viết này
Atomic function
Thư viện này dùng khá đơn giản và mục đích mà NVIDIA viết ra thư viện này để tránh trường hợp data hazard hay có thể nói là để đồng bộ các thread trong quá trình xử lí các phép toán
Ví dụ chúng ta có 1 đoạn code đơn giản
for(int i=0;i<n;i++){ y+=x[i]; }
Nhưng nếu đưa về song song
int tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid < n) { *d_result += d_data[tid]; }
Thì sẽ dẫn đến data hazard do không có sự đồng bộ giữa các thread và dẫn đến output sai, vậy nên atomic sinh ra để giải quyết vấn đề này
Có thể hiểu đơn giản atomic function hoạt động na ná giống for-loop , đều duyệt qua từng thread một ( thread nào khởi động trước sẽ được duyệt trước ) KHI VÀ CHỈ KHI các thread cùng read-write cùng 1 value còn nếu không thì vẫn song song như bình thường
Đây là code khi xử dụng atomic
int tid = threadIdx.x + blockDim.x * blockIdx.x; atomicAdd(result, array_add[tid]);
Và đây là 1 số atomic function
- 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 1 vài atomic function thú vị khác như atomicCAS, atomicExch
Full code mình sẽ để ở đây