CUDAには「Grid」、「Block」、「Thread」という単位がある。

「Thread」→いわゆる処理１つ分のこと。（例えばブラウザを起動しながらメールを起動するのは2つのスレッドとなるような感じ）
「Block」→複数個の「Thread」が集まったもの。
「Grid」→複数個の「Block」が集まったもの。

「ローカルメモリ」は「Thread」が使うメモリ。
「共有メモリ」は「Block」内で使うメモリ。つまりそのブロック内の「Thread」で共有される。
「グローバルメモリ」はGPU全体で使用できるメモリ。

CUDAではデータ配列を各々のスレッドが担当し、それぞれが並列に実行されるイメージ。
つまり、上の図の１スレッドに配列データの１要素分を割り当てる感じ。

CUDAの基本的な操作は次の通り。

1. GPUのメモリ確保
2. CPU→GPUへデータを転送
3. GPUで処理(__global__関数へ)
4. GPU→CPUでデータを取り出す

GPUにデータを転送するためには、まずGPU上にメモリを確保する必要がある。

【例】：GPU上にメモリ確保

const int AryNum = 20000;
…　…
//GPUのメモリ確保
int DataSize = sizeof(float)*AryNum;
float *gpuAry;
CUDA_SAFE_CALL(cudaMalloc( reinterpret_cast<void**>( &gpuAry), DataSize) );
…　…
cudaFree(gpuAry);//メモリ破棄

【解説】
cudaError_t cudaMalloc (void **devPtr, size_t size)
GPU上にメモリを確保する。
※CUDA_SAFE_CALL()でエラーを拾う。（自分でチェックしてもよい）
不要になったら
cudaError_t cudaFree (void *devPtr)
でメモリを破棄する。

GPUへのデータ転送はcudaMemcpy関数を使う
【例】：GPUへのデータ転送

//CPUでメモリ確保
float *ary = new float[AryNum];
…　値の設定　…

//GPUのメモリ確保
int DataSize = sizeof(float)*AryNum;
float *gpuAry;
…　GPUのメモリ確保　…

CUDA_SAFE_CALL( cudaMemcpy( gpuAry, ary, DataSize , cudaMemcpyHostToDevice));

【解説】
cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
「count」バイト分だけ「src」の内容を「dst」へコピー
「kind」には

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

が指定可能。

CPUでGPUが実行する関数をコールすることでGPUでの処理がはじまる。
※GPUでの関数の事をカーネルともよぶそうだ

【例】：カーネルの呼び出し

//------------- CUDAでの処理 ---------------//
__global__ void cuAddValue(float *gpuAryData)
{
	const int threadId = blockIdx.x * blockDim.x + threadIdx.x;  //スレッドIdを取得
	gpuAryData[threadId] += 20000;//すべての要素を+20000する

	__syncthreads(); //同期をとる
}

…　…
//GPU用にブロックとスレッドを用意
dim3 grid(100,1,1);//100×1サイズのグリッド(1グリッドあたり100ブロック)
dim3 block(200,1,1);//200×1のブロックサイズ(1ブロックあたり200スレッド)

//グリッド100個、ブロック200
cuAddValue<<<grid,block>>>(gpuAry);

詳細については

• CUDA編02::カーネルの呼び出し

「cudaMemcpy関数」の第４引数「にcudaMemcpyDeviceToHost」を指定すればよい。
【例】:GPUからデータを取り出す

//CPUでメモリ確保
float *ary = new float[AryNum];
…　値の設定　…

//GPUのメモリ確保
int DataSize = sizeof(float)*AryNum;
float *gpuAry;
…　GPUのメモリ確保　…

…　…
//GPUから戻す(GPU -> Host)
CUDA_SAFE_CALL( cudaMemcpy( ary,gpuAry ,DataSize, cudaMemcpyDeviceToHost) );

ここまでをまとめると次のようになる

const int AryNum = 20000;

//------------- CUDAでの処理 ---------------//
__global__ void cuAddValue(float *gpuAryData)
{
	…　…
}

//------------- メイン関数 --------------------//
int main(int argc, char **argv)
{
	CUT_DEVICE_INIT(argc,argv);

	//CPUでメモリ確保＋値設定
	float *ary = new float[AryNum];
	for(int loop = 0;loop < AryNum; ++loop)
	{
		ary[loop] = loop;
	}
	//GPUのメモリ確保
	int DataSize = sizeof(float)*AryNum;
	float *gpuAry;
	CUDA_SAFE_CALL(cudaMalloc( reinterpret_cast<void**>( &gpuAry), DataSize) );

	//GPUへデータをコピー(Host -> GPU)
	CUDA_SAFE_CALL( cudaMemcpy( gpuAry, ary, DataSize , cudaMemcpyHostToDevice));

	//GPU用にブロックとスレッドを用意
	dim3 grid(100,1,1);//100×1サイズのグリッド(1グリッドあたり100ブロック)
	dim3 block(200,1,1);//200×1のブロックサイズ(1ブロックあたり200スレッド)
	
	cuAddValue<<<grid,block>>>(gpuAry);

	//GPUから戻す(GPU -> Host)
	CUDA_SAFE_CALL( cudaMemcpy( ary,gpuAry ,DataSize, cudaMemcpyDeviceToHost) );
	cudaFree(gpuAry);//メモリ破棄

	//出力テスト
	for(int loop = 0;loop < AryNum; ++loop)
	{
		std::cout << ary[loop] << "\n";
	}

	delete [] ary;//メモリクリア

	CUT_EXIT(argc, argv);//終了

	return 0;
}

GPUへ渡した配列データを処理するにはなんとかして、アクセスするための添え字を取得する必要がある。
CUDAでは組み込み変数として、threadIdxやblockDimなどがある。
これらを使って、配列データにアクセスする。

詳細は

• CUDA編02::組み込み変数

【例】

//------------- CUDAでの処理 ---------------//
__global__ void cuAddValue(float *gpuAryData)
{
	const int threadId = blockIdx.x * blockDim.x + threadIdx.x;  //スレッドIdを取得
	gpuAryData[threadId] += 20000;//すべての要素を+20000する

	__syncthreads(); //同期をとる
}

…　…
cuAddValue<<<grid,block>>>(gpuAry);

【解説】
まず、現在のスレッド番号、ブロック番号から配列データのどこに対応するかを計算。
処理すべき配列番号を取得したらそこにアクセスして、何らかの処理をおこなう。

void __syncthreads()
ブロック内のスレッドがすべてその関数がかかれた所にくるまで待つ。
ブロック内で同期処理をする時に使う。

• CUDA編02::まとめ