CUDA编程 6 原子操作【参加CUDA线上训练营】

CUDA中的原子操作是指在多个线程并行执行时保证在某个特定时刻仅有一个线程能够访问共享资源（例如全局内存）的操作。

以下是常用的 CUDA 原子操作：

1. atomicAdd：对全局内存的指定位置进行加法操作，并返回该位置的旧值。

2. atomicSub：对全局内存的指定位置进行减法操作，并返回该位置的旧值。

3. atomicExch：对全局内存的指定位置执行交换操作，并返回该位置的旧值。

4. atomicMin：对全局内存的指定位置执行取最小值操作，并返回该位置的旧值。

5. atomicMax：对全局内存的指定位置执行取最大值操作，并返回该位置的旧值。

6. atomicAnd：对全局内存的指定位置执行位与操作，并返回该位置的旧值。

7. atomicOr：对全局内存的指定位置执行位或操作，并返回该位置的旧值。

8. atomicXor：对全局内存的指定位置执行位异或操作，并返回该位置的旧值。

9. atomicCAS：对全局内存的指定位置执行比较并交换操作，并返回该位置的旧值。

使用原子操作可以提高应用程序的并行性和效率。

课后作业：

1. 给定数组A[1000000]找出其中最大的值和最小的值

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <thrust/min_element.h>
#include <thrust/max_element.h>

__global__ void findMinMax(int *A, int *minValue, int *maxValue, int n)
{
    int tempMin = INT_MAX;
    int tempMax = INT_MIN;
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    while (i < n)
    {
        tempMin = min(tempMin, A[i]);
        tempMax = max(tempMax, A[i]);
        i += blockDim.x * gridDim.x;
    }

    __shared__ int sharedMin[1024];
    __shared__ int sharedMax[1024];
    sharedMin[threadIdx.x] = tempMin;
    sharedMax[threadIdx.x] = tempMax;
    __syncthreads();

    for (int j = blockDim.x / 2; j > 0; j /= 2)
    {
        if (threadIdx.x < j)
        {
            sharedMin[threadIdx.x] = min(sharedMin[threadIdx.x], sharedMin[threadIdx.x + j]);
            sharedMax[threadIdx.x] = max(sharedMax[threadIdx.x], sharedMax[threadIdx.x + j]);
        }
        __syncthreads();
    }

    if (threadIdx.x == 0)
    {
        *minValue = sharedMin[0];
        *maxValue = sharedMax[0];
    }
}

int main()
{
    int n = 1000000;
    int *A, *minValue, *maxValue;

    cudaMalloc((void **)&A, n * sizeof(int));
    cudaMalloc((void **)&minValue, sizeof(int));
    cudaMalloc((void **)&maxValue, sizeof(int));

    // fill A with data

    findMinMax<<<(n + 1023) / 1024, 1024>>>(A, minValue, maxValue, n);

    int h_minValue, h_maxValue;
    cudaMemcpy(&h_minValue, minValue, sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(&h_maxValue, maxValue, sizeof(int), cudaMemcpyDeviceToHost);

    printf("Min value: %d\n", h_minValue);
    printf("Max value: %d\n", h_maxValue);

    cudaFree(A);
    cudaFree(minValue);
    cudaFree(maxValue);

    return 0;
}

2. 给定数组A[1000000]找出其中最大的十个值

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>

__global__ void findTopTen(int *A, int *topTen, int n)
{
    int temp[10];
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    for (int j = 0; j < 10; j++)
    {
        temp[j] = INT_MIN;
    }

    while (i < n)
    {
        for (int j = 0; j < 10; j++)
        {
            if (A[i] > temp[j])
            {
                for (int k = 9; k > j; k--)
                {
                    temp[k] = temp[k - 1];
                }
                temp[j] = A[i];
                break;
            }
        }
        i += blockDim.x * gridDim.x;
    }

    __shared__ int shared[1024][10];
    for (int j = 0; j < 10; j++)
    {
        shared[threadIdx.x][j] = temp[j];
    }
    __syncthreads();

    for (int j = blockDim.x / 2; j > 0; j /= 2)
    {
        for (int k = 0; k < 10; k++)
        {
            if (threadIdx.x < j)
            {
                if (shared[threadIdx.x + j][k] > shared[threadIdx.x][k])
                {
                    shared[threadIdx.x][k] = shared[threadIdx.x + j][k];
                }
            }
            __syncthreads();
        }
    }

    if (threadIdx.x == 0)
    {
        for (int j = 0; j < 10; j++)
        {
            topTen[j] = shared[0][j];
        }
    }
}

int main()
{
    int n = 1000000;
    int *A, *topTen;

    cudaMalloc((void **)&A, n * sizeof(int));
    cudaMalloc((void **)&topTen, 10 * sizeof(int));

    // fill A with data

    findTopTen<<<(n + 1023) / 1024, 1024>>>(A, topTen, n);

    int h_topTen[10];
    cudaMemcpy(h_topTen, topTen, 10 * sizeof(int), cudaMemcpyDeviceToHost);

    for (int i = 0; i < 10; i++)
    {
        printf("%d ", h_topTen[i]);
    }
    printf("\n");

    cudaFree(A);
    cudaFree(topTen);

    return 0;
}