CUDA的原子操作

CUDA的原子操作是用于在并行计算中保证多个线程对共享数据进行安全访问的一种机制。由于CUDA允许多个线程同时执行，这可能会导致数据竞争和不一致性，原子操作通过确保某个操作在执行过程中不会被其他线程打断，从而解决这些问题。

原子操作的特点

1. 原子性：

   • 原子操作是不可分割的操作。在执行原子操作时，其他线程无法对该数据进行读写，确保了数据的一致性。

2. 支持的数据类型：

   • CUDA提供了对基本数据类型（如整数、浮点数）的原子操作支持，包括 int, unsigned int, float 等。

3. 常用的原子操作：

   • atomicAdd(): 对变量进行加法操作。
   • atomicSub(): 对变量进行减法操作。
   • atomicExch(): 交换两个变量的值。
   • atomicMin(): 更新变量为最小值。
   • atomicMax(): 更新变量为最大值。
   • atomicCAS(): 原子比较并交换。

#include <stdio.h>
#include <cuda.h>

__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
    // 计算当前线程的全局索引
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 确保索引不越界
    if (i < N) {
        C[i] = A[i] + B[i]; // 执行向量加法
    }
}

__global__ void vectorSum(const float *C, float *result, int N) {
    // 计算当前线程的全局索引
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 使用原子操作累加结果
    if (i < N) {
        atomicAdd(result, C[i]); // 将C[i]的值累加到result中
    }
}

int main() {
    int N = 1 << 5; // 向量大小
    printf("N = %d\n", N);
    size_t size = N * sizeof(float);

    // 在主机上分配内存
    float *h_A = (float *)malloc(size);
    float *h_B = (float *)malloc(size);
    float *h_C = (float *)malloc(size);
    float h_result = 0.0f; // 用于存储结果

    // 初始化向量
    for (int i = 0; i < N; i++) {
        h_A[i] = i;
        h_B[i] = i * 2.0f;
    }

    // 在设备上分配内存
    float *d_A, *d_B, *d_C, *d_result;
    cudaMalloc((void **)&d_A, size);
    cudaMalloc((void **)&d_B, size);
    cudaMalloc((void **)&d_C, size);
    cudaMalloc((void **)&d_result, sizeof(float));

    // 将数据从主机复制到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_result, &h_result, sizeof(float), cudaMemcpyHostToDevice);

    // 定义线程块大小和网格大小
    int threadsPerBlock = 256; // 每个线程块中的线程数
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    printf("blocksPerGrid = %d\n", blocksPerGrid);

    // 启动内核进行向量加法
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 启动内核进行结果求和
    vectorSum<<<blocksPerGrid, threadsPerBlock>>>(d_C, d_result, N);

    // 将结果从设备复制回主机
    cudaMemcpy(&h_result, d_result, sizeof(float), cudaMemcpyDeviceToHost);

    printf("Sum of vector C: %f\n", h_result); // 输出C向量的和

    // 释放内存
    free(h_A);
    free(h_B);
    free(h_C);
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    cudaFree(d_result);

    printf("Completed successfully!\n");
    return 0;
}