Ở 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