CUDA共有メモリ操作(_shared__キーワード)

CUDAによるGPUのメモリ操作の場合、スレッドがあるデータに対して頻繁に読み書き操作を行うと、操作するデータ常駐キャッシュを設定することができ、コードの実行効率をさらに向上させ、同じスレッドブロック内のすべてのスレッドがそのメモリ領域を共有することができる.しかし、複数のスレッドが同じメモリ領域を操作する場合、競合を回避するためにスレッドを同期する必要があります.Cudaで使用_shared__キーワード、ここでは__syncthreads()はスレッドの同期を制御します.
コア関数dotに示すように、コードは配列を生成し、配列内のすべての要素の和を計算し、集計(Reduction)メソッドを使用します.
すべての開始スレッドが最初に_syncthreads()の前のすべてのコードが実行されると、すべてのスレッドが実行された後のコードを続行します.そうしないと、cache配列が初期化されていないうちにすべての要素の計算を開始する可能性があります.このように計算結果は明らかに間違っています.

#include 

#define min(a, b) (a>>(dev_a, dev_b, dev_partial_c);

	//    dev_partial_c   GPU     CPU
	cudaMemcpy(partial_c, dev_partial_c, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost);

	// CPU          
	c = 0.0;
	for(int i = 0; i < blocksPerGrid; i++){
		c += partial_c[i];
	}

	printf("%s
", "======================================");
	for(int i = 0; i < N; i++){
		printf("%f  %f 
", a[i], b[i]);
	}

	printf("c = %f 
", c);
	printf("blocksPerGrid %d 
", blocksPerGrid);

	//    GPU     
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_partial_c);

	//    CPU     
	free(a);
	free(b);
	free(partial_c);

	return 0;
}