CUDAについてメモ

各ブロックでどう同期するか

__synchthreadsはブロック内のスレッドを同期はするが、ブロック同士では同期しない。いったんcuda kernelを終了させれば全ブロックで同期される。しかしこの方法の場合、kernelで確保したshared memoryはいったん破棄しなくてはならない。kernelの外側でshared memoryの宣言を書いてもやはり一旦破棄されるようだ。
kernelを終了せずに同期できないだろか?
VRAMのある領域をカウンタに使って各ブロックでインクリメントさせて、全ブロック数になったら同期されたとして次に進むという方法でいいのだろうか?この方法がどれくらいコストがかかるのか調べる必要があるかもしれない。ちなみにnVidiaは全ブロックでの同期を推奨しないらしい。なぜかは知らない。

二つのデバイスをどう扱うか

研究室のマシンにはnVidiaのグラフィックボードが二つ刺さっているのだが、SLIが有効でないためにCUDAでは別々のデバイスとしてみえてしまう。
それをどうやって扱うかについて調べた。

このサイト

を参考に最小限必要なコードをつくってみた。

#include <stdio.h>
#include <cutil.h>
#include <multithreading.h>

#define MAX_CPU_THREAD 2

__global__ void gpuKernel(int dev, int* gpu) 
{
	*gpu = dev;
}

// ホストメモリ
int data[MAX_CPU_THREAD];

//GPUの数だけ下記の関数が実行される
CUT_THREADPROC gpuThread(int* pdev)
{
	int& dev = *pdev;
	// デバイスの指定
	CUDA_SAFE_CALL(cudaSetDevice(dev));

	int* gpu;
	// グローバルメモリの確保
	CUDA_SAFE_CALL(cudaMalloc((void**)&gpu, sizeof(int)));
	// カーネルの呼び出し
	gpuKernel<<< 1, 1 >>>(dev,gpu);
	// グローバルメモリからホストのメモリに転送
	CUDA_SAFE_CALL(cudaMemcpy(data+dev, gpu, sizeof(int), cudaMemcpyDeviceToHost)); 
	// グローバルメモリの開放
	CUDA_SAFE_CALL(cudaFree(gpu)); 

	// スレッドの終わり(不要?)
	CUT_THREADEND;
}


int main( int argc, char** argv) 
{

	// GPUの数
	int s_gpuCount;
	CUDA_SAFE_CALL(cudaGetDeviceCount(&s_gpuCount));
	printf("%d GPUs found\n", s_gpuCount);

	// ThreadのID
	int threadIds[MAX_CPU_THREAD];
	// Thread情報
	CUTThread* threads = (CUTThread *)malloc(sizeof(CUTThread) * s_gpuCount);

	//デバイスの数
	for(int i = 0; i < s_gpuCount; i++)
	{
		threadIds[i] = i;
		// スレッド呼び出し
		threads[i] = cutStartThread((CUT_THREADROUTINE)gpuThread, (void *)&threadIds[i]);
	}

	//すべてのスレッドが終了するのを待つ
	cutWaitForThreads(threads, s_gpuCount);
	for(int i = 0; i < s_gpuCount; i++)
	{
		printf("%d\n", data[i]);
	}
	free(threads);
	return 0;
}

各デバイス向けにCPU側で二つスレッドを立てているもよう。

ちなみにcudaMemcpyDeviceToDeviceで各デバイスで確保したVRAMのアドレスを指定してコピーさせようとしたがうまくいかなかった。
両方共ポインタとしては同じアドレスになっていて、cudaMemcpyをしても同じものがコピーされるだけだった。