ストリーム

非同期同時実行†

ホストとデバイスの同時実行を簡単にするために， 以下の関数は非同期(デバイスが要求されたタスクを完了する前にホストスレッドに制御が戻る)になっています．

• カーネル起動関数
• 接尾辞がAsyncのメモリコピー処理関数
• デバイス-デバイス間メモリコピー処理関数
• メモリ設定関数

ページロックホストメモリでも述べたように，デバイスによっては Page-lockedホストメモリとデバイスメモリ間のコピーもカーネル実行と同時に行えます． これらの非同期関数を管理するためにストリーム，モニターするためにイベントという機能が用意されています．

ストリーム†

ストリームは順に実行される命令列で，非同期の処理の流れを示します． ストリームを作成することで，メモリの確保，データ転送，カーネル実行，データを戻す，メモリの解放 の基本的なGPUの処理の流れを並列に実行できます．

ストリームはストリームオブジェクトを作成して， 各非同期関数の引数にそれを指定することでその中で実行する命令を指定します． 以下にストリーム作成と命令の指定例として，1からmax_numまでの数の総和計算をnstream個のストリームで 実行するコードです．ただし，max_numはnstreamの倍数である必要があります． カーネルにはCUDA SDKに同梱のサンプルコードReductionのreduce2カーネル関数を用いています．

void CuStreamTest(int max_num)
{
	int nstream = 2;

	int n = max_num/nstream;
	int size = n*sizeof(int);

	int **dInData = (int**)malloc(nstream*sizeof(int*));
	int **hInData = (int**)malloc(nstream*sizeof(int*));
	for(int i = 0; i < nstream; ++i){
		cutilSafeCall(cudaMallocHost((void**)&hInData[i], size));
		cutilSafeCall(cudaMalloc((void**)&dInData[i], size));

		for(int j = 0; j < n; ++j){
			hInData[i][j] = i*n+j;
		}
	}

	uint numThreads, numBlocks;
	numThreads = min(64, n);
	numBlocks = iDivUp(n, numThreads);
    int smemSize = numThreads*sizeof(int);

	int **hOutData = (int**)malloc(nstream*sizeof(int*));
	int **dOutData = (int**)malloc(nstream*sizeof(int*));
	for(int i = 0; i < nstream; ++i){
		cutilSafeCall(cudaMallocHost((void**)&hOutData[i], numBlocks*sizeof(int)));
		cutilSafeCall(cudaMalloc((void**)&dOutData[i], numBlocks*sizeof(int)));
	}

	cudaStream_t *streams = (cudaStream_t*)malloc(nstream*sizeof(cudaStream_t));
	for(int i = 0; i < nstream; ++i){
		cutilSafeCall(cudaStreamCreate(&streams[i]));
	}

	for(int i = 0; i < nstream; ++i){
	    cutilSafeCall(cudaMemcpyAsync(dInData[i], hInData[i], size, cudaMemcpyHostToDevice, streams[i]));
	}
	
	for(int i = 0; i < nstream; ++i){
		reduce2<int><<< numBlocks, numThreads, smemSize, streams[i] >>>(dInData[i], dOutData[i], n);
	}

	for(int i = 0; i < nstream; ++i){
	    cutilSafeCall(cudaMemcpyAsync(hOutData[i], dOutData[i], numBlocks*sizeof(int), cudaMemcpyDeviceToHost, streams[i]));
	}

	cudaThreadSynchronize();

	int sum_gpu = 0;
	for(int i = 0; i < nstream; ++i){
		for(int j = 0; j < numBlocks; ++j){
			sum_gpu += hOutData[i][j];
		}
	}

	int sum_cpu = 0;
	for(int j = 0; j < n*nstream; ++j){
		sum_cpu += j;
	}

	printf("gpu : %d\n", sum_gpu);
	printf("cpu : %d\n", sum_cpu);

	for(int i = 0; i < nstream; i++){
        cudaStreamDestroy(streams[i]);
	}

	for(int i = 0; i < nstream; i++){
		cudaFreeHost(hInData[i]);
		cudaFreeHost(hOutData[i]);
		cudaFree(dInData[i]);
		cudaFree(dOutData[i]);
	}

    cudaThreadExit();
}

実行結果は，

gpu : 8256
cpu : 8256

です．

イベント†