CUDA高性能并行计算(9)--原子函数及其合理使用

CUDA全局内存、共享内存的合理使用

原子操作

所谓原子操作，就是不能被中断的操作，因为在多线程任务中，假设有全局内存可以被所有线程访问，且线程的任务均为读--将值加一--写回，那么当线程1进行加一的操作时，线程2将其结果写入到该内存中，那么接下来线程1再写回时，就覆盖了线程2的计算结果。如果这个任务被用于计数的话，那么结果就是线程之间串扰导致错误。

因此，我们希望某个线程在执行这一套流程时，其他线程不能访问这个全局内存，只有这样我们才能保证最终结果的正确性，这就是原子操作。

当然，假如说每个线程执行的任务都是需要对共享内存进行大量的原子读写，就会导致性能下降的很多。

我们可以使用形如atomicxxx(addr, val)的函数来进行原子操作，**原子函数都是__device__修饰的**。

函数	功能
T atomicAdd(T*, T)	加法
T atomicSub(T*, T)	减法
T atomicExch(T*, T)	交换
T atomicMin(T*, T)	最小值
T atomicMax(T*, T)	最大值
T atomicInc(T*, T)	自增
T atomicDec(T*, T)	自减
T atomicCAS(T*, T1, T2)	比较-交换
T atomicAnd(T*, T)	按位与
T atomicOr(T*, T)	按位或
T atomicXor(T*, T)	按位异或

其结果一般都是 *(T*) = *(T *) operate T

原子函数的返回值都是*(T *)的原始值。

举例：T atomicSub(T* addr, T val)

等效于

T tmp = *addr;
*addr = *addr - val;
return tmp;

其中比较特殊的T atomicCAS(T*, T1, T2)的作用是*(T*) = (*(T*) = T1) ? T2 : *(T*)

还有需要注意一点，原子函数是有精度限制的，比如atomicSub要求数据类型为int,unsigned,unsigned long long，而不能是浮点数类型。

一般所有的原子操作都是通过atomicCAS来实现，但是在实际过程中要避免这样做，因为性能会下降。

下面给出一个例子

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_atomic_functions.h"
#include <stdio.h>
#include <stdlib.h>

#define ATOMIC 0
#define TPB 32
#define N 1024


__global__ void dotKernel(int * d_res, const int *d_a, const int * d_b, int n)
{
	const int idx = threadIdx.x + blockDim.x * blockIdx.x;
	const int s_idx = threadIdx.x;

	if (idx >= n)
		return;

	__shared__ int s_prod[TPB];
	s_prod[s_idx] = d_b[idx] * d_a[idx];
	__syncthreads(); // 至此，所有的分块计算都完成了

	if (s_idx == 0) // 使用每个线程块中的第一个线程进行数据合并计算
	{
		int blockSum = 0;
		for (int j = 0; j < blockDim.x; ++j)
		{
			blockSum += s_prod[j];
		}

		printf("block %d, blocksum %d\n", blockIdx.x, blockSum);

#if ATOMIC == 1
		atomicAdd(d_res, blockSum); // 原子计算
#else
		*d_res += blockSum; // 直接计算
#endif
	}
}

void doLauncher(int * res, const int *a, const int *b, int n)
{
	int * d_res;
	int *d_a = 0;
	int *d_b = 0;

	// 分配内存
	cudaMalloc(&d_res, sizeof(int));
	cudaMalloc(&d_a, n * sizeof(int));
	cudaMalloc(&d_b, n * sizeof(int));

	// 拷贝数据
	cudaMemset(d_res, 0, sizeof(int));
	cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
	
	// 调用核函数计算并将数据写回主机端
	dotKernel << <(n + TPB - 1) / TPB, TPB >> > (d_res, d_a, d_b, n);
	cudaMemcpy(res, d_res, sizeof(int), cudaMemcpyDeviceToHost);

	// 释放内存
	cudaFree(d_res);
	cudaFree(d_a);
	cudaFree(d_b);
}

int main()
{
	int cpu_res = 0;
	int gpu_res = 0;
	int *a = (int *)malloc(N * sizeof(int));
	int *b = (int *)malloc(N * sizeof(int));

	for (int i = 0; i < N; ++i)
	{
		a[i] = 1;
		b[i] = 1;
	}

	// 先调用CPU计算结果
	for (int i = 0; i < N; i++)
	{
		cpu_res += a[i] * b[i];
	}

	printf("\n\ncpu res = %d\n\n", cpu_res);

	// 调用GPU对比计算结果
	doLauncher(&gpu_res, a, b, N);

	printf("\n\ngpu res = %d\n\n", gpu_res);

	free(a);
	free(b);
	return 0;
}

可以修改宏ATOMIC来查看使用原子操作和非原子操作的结果。其中原子操作每次都能给出正确结果，而非原子操作每次结果都不一样，且都不正确。

总结

• 当涉及归约或者其他类型的需要多个线程访问同一个地址时，需要使用原子操作。