CUDAアトミック操作

<スパン CUDAのアトミック演算は、変数に対する"read-modify-write"の3つの演算の実行の最小単位と理解でき、より小さな部分に分解できず、その実行中に他の並列スレッドが変数に対して読み書きをすることを許さないものである。この仕組みに基づき、アトミック演算は複数のスレッドで共有される変数の相互排他的保護を実現し、変数に対するいかなる単一演算の結果も正しいことを保証する。

アトミック操作により、複数の並列スレッドで共有されるメモリの読み書きを保護し、一度に1つのスレッドのみが変数への読み書きを行うことができます。 あるスレッドが変数に対して操作を行うと、同様に操作を行いたい他のスレッドは、前のスレッドの実行が完了するまで待たなければならない。アトミック操作では、パフォーマンスを犠牲にして安全性を確保します。

<スパン

CUDAは様々なアトミック操作をサポートしていますが、一般的なものは以下の通りです。

<スパン

1. atomicAdd()

intomicAdd(int* address, int val);
unsigned intomicAdd(unsigned int* address, unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,unsigned long long int val);
グローバルメモリまたは共有メモリのアドレスにある32ビットまたは64ビットワードoldを読み出し、（old + val）を計算し、メモリ内の同じアドレスに結果を格納します。この3つの操作は、1つのアトミックトランザクションで実行されます。グローバルメモリのみが64ビットワードをサポートします。

2. atomicSub()

intomicSub(int* address, int val);
unsigned intomicSub(unsigned int* address, unsigned int val);
グローバルメモリまたは共有メモリのアドレスにある32ビットワードoldを読み出し、(old - val)を計算し、結果をメモリの同じアドレスに格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数は old を返します。
3.アトミックエクスチェンジ()

atomicExch(int* address, int val);
unsigned int atomicExch(unsigned int* address, unsigned int val);
unsigned long long int atomicExch(unsigned long long int* address,unsigned long long int val);
float atomicExch(float* address, float val);

グローバルメモリまたは共有メモリのアドレスaddressにある32ビットまたは64ビットのワードoldを読み出し、メモリの同じアドレスにvalを格納します。この2つの操作は、1つのアトミックトランザクションで実行されます。グローバルメモリのみが64ビットワードをサポートします。
4. atomicMin()

intomicMin(int* address, int val);
unsigned intomicMin(unsigned int* address, unsigned int val);

グローバルメモリまたは共有メモリのアドレスaddressにある32ビットワードoldを読み出し、oldとvalの最小値を計算し、結果をメモリの同じアドレスに格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数はoldを返します。
5.アトミックマックス()

intomicMax(int* address, int val);
unsigned intomicMax(unsigned int* address, unsigned int val);

グローバルメモリまたは共有メモリのアドレスaddressにある32ビットワードoldを読み出し、oldとvalの最大値を計算し、結果をメモリの同じアドレスに格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数はoldを返します。
6. atomicInc()

unsigned atomicInc(unsigned int* address,unsigned int val);

グローバルメモリまたは共有メモリのAddressにある32ビットワードoldを読み出し、((old >= val) ?を計算します。0 : (old+1)) を計算し、その結果を同じアドレスのメモリに格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数はoldを返します．

7. atomicDec()

unsigned intomicDec(unsigned int* address,unsigned int val);

グローバルメモリまたは共有メモリのアドレスにある32ビットワードoldを読み込み、(((old == 0) | (old > val)) ? val : (old-1)) を計算し、同じアドレスのメモリに結果を格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数はoldを返します．

8. atomicCAS()

intomicCAS(int* address, int compare, int val);
unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address,unsigned long long int compare,unsigned long long int val);

グローバルメモリまたは共有メモリのアドレスaddressにある32ビットまたは64ビットワードoldを読み込み、(old == compare ? val : old)を計算し、同じアドレスのメモリに結果を格納します。これら 3 つの操作は、1 つのアトミック トランザクションで実行されます。この関数は old (compare と swap) を返します。グローバルメモリのみ64ビットワードをサポートします。
9. atomicAnd()

intomicAnd(int* address, int val);
unsigned intomicAnd(unsigned int* address, unsigned int val);

グローバルメモリまたは共有メモリのアドレスにある32ビットワードoldを読み出し、(old & val)を計算し、結果をメモリの同じアドレスに格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数は old を返します。
10. atomicOr()

intomicOr(int* address, int val);
unsigned intomicOr(unsigned int* address, unsigned int val);

グローバルメモリまたは共有メモリのアドレスにある32ビットワードoldを読み出し、(old | val)を計算し、結果をメモリの同じアドレスに格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数はoldを返します。

11. atomicXor()

intomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address, unsigned int val);

グローバルメモリまたは共有メモリのアドレスにある32ビットワードoldを読み出し、(old ^ val)を計算し、結果をメモリの同じアドレスに格納します。これら3つの操作は、1つのアトミックトランザクションで実行されます。この関数はoldを返します。

例として、1024個のスレッドを定義し、その1024個のスレッドのIDの合計を求め、それぞれがsum変数sumにアクセスし、アトミック演算をしないと、実行結果がおかしくなり不定になります。

#include <stdio.h>    
#include <stdlib.h>   
#include <cuda_runtime.h>  

#define SIZE 1024

__global__ void histo_kernel(int size, unsigned int *histo)
{
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	if (i < size)
	{
		//*histo+=i;
		atomicAdd(histo, i);
	}
}

int main(void)
{
	int threadSum = 0;

	//allocate memory and copy initial data
	unsigned int *dev_histo;

	cudaMalloc((void**)&dev_histo, sizeof(int));
	cudaMemcpy(dev_histo, &threadSum, sizeof(int), cudaMemcpyHostToDevice);

	// kernel launch - 2x the number of mps gave best timing  
	cudaDeviceProp prop;
	cudaGetDeviceProperties(&prop, 0);

	int blocks = prop.multiProcessorCount;
	// make sure there are enough threads
	histo_kernel << <blocks * 2, (SIZE + 2 * blocks - 1) / blocks / 2 >> > (SIZE, dev_histo);

	// copy data back to CPU memory
	cudaMemcpy(&threadSum, dev_histo, sizeof(int), cudaMemcpyDeviceToHost);
	printf("Threads SUM: %d\n", threadSum);
	getchar();
	cudaFree(dev_histo);
	return 0;
}