次のコードは、配列内のすべての要素を各要素グループ32の最初の要素に合計します。32
int i = threadIdx.x;
int warpid = i&31;
if(warpid < 16){
s_buf[i] += s_buf[i+16];__syncthreads();
s_buf[i] += s_buf[i+8];__syncthreads();
s_buf[i] += s_buf[i+4];__syncthreads();
s_buf[i] += s_buf[i+2];__syncthreads();
s_buf[i] += s_buf[i+1];__syncthreads();
}
__syncthreads()すべての操作が同じワープで行われるため、コード内のすべてを削除できると思いました。しかし、それらを削除すると、ガベージの結果が返されます。パフォーマンスにはあまり影響しませんが、なぜ__syncthreads()ここが必要なのか知りたいです。
上記の2つは完全に満足できるものではないと思うので、ここで回答します。この回答の「知的財産」は、このプレゼンテーション(スライド 22) でこの問題を指摘した Mark Harris と、上記のコメントで OP にこの問題を指摘した @talonmies に属します。
最初に、OPが求めていたものを再開して、彼の間違いをフィルタリングしてみましょう.
OPは、共有メモリの削減、ループ展開によるワープ削減の最後のステップを扱っているようです。彼は次のようなことをしています
template <class T>
__device__ void warpReduce(T *sdata, int tid) {
sdata[tid] += sdata[tid + 32];
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid + 8];
sdata[tid] += sdata[tid + 4];
sdata[tid] += sdata[tid + 2];
sdata[tid] += sdata[tid + 1];
}
template <class T>
__global__ void reduce4_no_synchthreads(T *g_idata, T *g_odata, unsigned int N)
{
extern __shared__ T sdata[];
unsigned int tid = threadIdx.x; // Local thread index
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; // Global thread index - Fictitiously double the block dimension
// --- Performs the first level of reduction in registers when reading from global memory.
T mySum = (i < N) ? g_idata[i] : 0;
if (i + blockDim.x < N) mySum += g_idata[i+blockDim.x];
sdata[tid] = mySum;
// --- Before going further, we have to make sure that all the shared memory loads have been completed
__syncthreads();
// --- Reduction in shared memory. Only half of the threads contribute to reduction.
for (unsigned int s=blockDim.x/2; s>32; s>>=1)
{
if (tid < s) { sdata[tid] = mySum = mySum + sdata[tid + s]; }
// --- At the end of each iteration loop, we have to make sure that all memory operations have been completed
__syncthreads();
}
// --- Single warp reduction by loop unrolling. Assuming blockDim.x >64
if (tid < 32) warpReduce(sdata, tid);
// --- Write result for this block to global memory. At the end of the kernel, global memory will contain the results for the summations of
// individual blocks
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
Mark Harris と talonmies が指摘したように、コンパイラの最適化を防ぐには、共有メモリ変数sdataを として宣言する必要があります。volatileしたがって、__device__上記の関数を定義する正しい方法は次のとおりです。
template <class T>
__device__ void warpReduce(volatile T *sdata, int tid) {
sdata[tid] += sdata[tid + 32];
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid + 8];
sdata[tid] += sdata[tid + 4];
sdata[tid] += sdata[tid + 2];
sdata[tid] += sdata[tid + 1];
}
上記の 2 つのケースに対応する逆アセンブルされたコード、つまり、 not orsdataとして宣言されたコード (Fermi アーキテクチャ用にコンパイルされたコード) を見てみましょう。volatilevolatile
いいえvolatile
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */
/*0010*/ SHL R3, R0, 0x1; /* 0x6000c0000400dc03 */
/*0018*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */
/*0020*/ IMAD R3, R3, c[0x0][0x8], R2; /* 0x200440002030dca3 */
/*0028*/ IADD R4, R3, c[0x0][0x8]; /* 0x4800400020311c03 */
/*0030*/ ISETP.LT.U32.AND P0, PT, R3, c[0x0][0x28], PT; /* 0x188e4000a031dc03 */
/*0038*/ ISETP.GE.U32.AND P1, PT, R4, c[0x0][0x28], PT; /* 0x1b0e4000a043dc03 */
/*0040*/ @P0 ISCADD R3, R3, c[0x0][0x20], 0x2; /* 0x400040008030c043 */
/*0048*/ @!P1 ISCADD R4, R4, c[0x0][0x20], 0x2; /* 0x4000400080412443 */
/*0050*/ @!P0 MOV R5, RZ; /* 0x28000000fc0161e4 */
/*0058*/ @!P1 LD R4, [R4]; /* 0x8000000000412485 */
/*0060*/ @P0 LD R5, [R3]; /* 0x8000000000314085 */
/*0068*/ SHL R3, R2, 0x2; /* 0x6000c0000820dc03 */
/*0070*/ NOP; /* 0x4000000000001de4 */
/*0078*/ @!P1 IADD R5, R4, R5; /* 0x4800000014416403 */
/*0080*/ MOV R4, c[0x0][0x8]; /* 0x2800400020011de4 */
/*0088*/ STS [R3], R5; /* 0xc900000000315c85 */
/*0090*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0098*/ MOV R6, c[0x0][0x8]; /* 0x2800400020019de4 */
/*00a0*/ ISETP.LT.U32.AND P0, PT, R6, 0x42, PT; /* 0x188ec0010861dc03 */
/*00a8*/ @P0 BRA 0x118; /* 0x40000001a00001e7 */
/*00b0*/ NOP; /* 0x4000000000001de4 */
/*00b8*/ NOP; /* 0x4000000000001de4 */
/*00c0*/ MOV R6, R4; /* 0x2800000010019de4 */
/*00c8*/ SHR.U32 R4, R4, 0x1; /* 0x5800c00004411c03 */
/*00d0*/ ISETP.GE.U32.AND P0, PT, R2, R4, PT; /* 0x1b0e00001021dc03 */
/*00d8*/ @!P0 IADD R7, R4, R2; /* 0x480000000841e003 */
/*00e0*/ @!P0 SHL R7, R7, 0x2; /* 0x6000c0000871e003 */
/*00e8*/ @!P0 LDS R7, [R7]; /* 0xc10000000071e085 */
/*00f0*/ @!P0 IADD R5, R7, R5; /* 0x4800000014716003 */
/*00f8*/ @!P0 STS [R3], R5; /* 0xc900000000316085 */
/*0100*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0108*/ ISETP.GT.U32.AND P0, PT, R6, 0x83, PT; /* 0x1a0ec0020c61dc03 */
/*0110*/ @P0 BRA 0xc0; /* 0x4003fffea00001e7 */
/*0118*/ ISETP.GT.U32.AND P0, PT, R2, 0x1f, PT; /* 0x1a0ec0007c21dc03 */
/*0120*/ @P0 BRA.U 0x198; /* 0x40000001c00081e7 */
/*0128*/ @!P0 LDS R8, [R3]; /* 0xc100000000322085 */
/*0130*/ @!P0 LDS R5, [R3+0x80]; /* 0xc100000200316085 */
/*0138*/ @!P0 LDS R4, [R3+0x40]; /* 0xc100000100312085 */
/*0140*/ @!P0 LDS R7, [R3+0x20]; /* 0xc10000008031e085 */
/*0148*/ @!P0 LDS R6, [R3+0x10]; /* 0xc10000004031a085 */
/*0150*/ @!P0 IADD R8, R8, R5; /* 0x4800000014822003 */
/*0158*/ @!P0 IADD R8, R8, R4; /* 0x4800000010822003 */
/*0160*/ @!P0 LDS R5, [R3+0x8]; /* 0xc100000020316085 */
/*0168*/ @!P0 IADD R7, R8, R7; /* 0x480000001c81e003 */
/*0170*/ @!P0 LDS R4, [R3+0x4]; /* 0xc100000010312085 */
/*0178*/ @!P0 IADD R6, R7, R6; /* 0x480000001871a003 */
/*0180*/ @!P0 IADD R5, R6, R5; /* 0x4800000014616003 */
/*0188*/ @!P0 IADD R4, R5, R4; /* 0x4800000010512003 */
/*0190*/ @!P0 STS [R3], R4; /* 0xc900000000312085 */
/*0198*/ ISETP.NE.AND P0, PT, R2, RZ, PT; /* 0x1a8e0000fc21dc23 */
/*01a0*/ @P0 BRA.U 0x1c0; /* 0x40000000600081e7 */
/*01a8*/ @!P0 ISCADD R0, R0, c[0x0][0x24], 0x2; /* 0x4000400090002043 */
/*01b0*/ @!P0 LDS R2, [RZ]; /* 0xc100000003f0a085 */
/*01b8*/ @!P0 ST [R0], R2; /* 0x900000000000a085 */
/*01c0*/ EXIT; /* 0x8000000000001de7 */
行/*0128*/-/*0148*/、/*0160*/および/*0170*/は、レジスターへの共有メモリーのロードと、レジスター/*0190*/からの共有メモリーストアへの行に対応します。中間行は、レジスタで実行される合計に対応します。そのため、中間結果はレジスター (各スレッド専用) に保持され、毎回共有メモリーにフラッシュされることはなく、スレッドが中間結果を完全に可視化できなくなります。
volatile
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */
/*0010*/ SHL R3, R0, 0x1; /* 0x6000c0000400dc03 */
/*0018*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */
/*0020*/ IMAD R3, R3, c[0x0][0x8], R2; /* 0x200440002030dca3 */
/*0028*/ IADD R4, R3, c[0x0][0x8]; /* 0x4800400020311c03 */
/*0030*/ ISETP.LT.U32.AND P0, PT, R3, c[0x0][0x28], PT; /* 0x188e4000a031dc03 */
/*0038*/ ISETP.GE.U32.AND P1, PT, R4, c[0x0][0x28], PT; /* 0x1b0e4000a043dc03 */
/*0040*/ @P0 ISCADD R3, R3, c[0x0][0x20], 0x2; /* 0x400040008030c043 */
/*0048*/ @!P1 ISCADD R4, R4, c[0x0][0x20], 0x2; /* 0x4000400080412443 */
/*0050*/ @!P0 MOV R5, RZ; /* 0x28000000fc0161e4 */
/*0058*/ @!P1 LD R4, [R4]; /* 0x8000000000412485 */
/*0060*/ @P0 LD R5, [R3]; /* 0x8000000000314085 */
/*0068*/ SHL R3, R2, 0x2; /* 0x6000c0000820dc03 */
/*0070*/ NOP; /* 0x4000000000001de4 */
/*0078*/ @!P1 IADD R5, R4, R5; /* 0x4800000014416403 */
/*0080*/ MOV R4, c[0x0][0x8]; /* 0x2800400020011de4 */
/*0088*/ STS [R3], R5; /* 0xc900000000315c85 */
/*0090*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0098*/ MOV R6, c[0x0][0x8]; /* 0x2800400020019de4 */
/*00a0*/ ISETP.LT.U32.AND P0, PT, R6, 0x42, PT; /* 0x188ec0010861dc03 */
/*00a8*/ @P0 BRA 0x118; /* 0x40000001a00001e7 */
/*00b0*/ NOP; /* 0x4000000000001de4 */
/*00b8*/ NOP; /* 0x4000000000001de4 */
/*00c0*/ MOV R6, R4; /* 0x2800000010019de4 */
/*00c8*/ SHR.U32 R4, R4, 0x1; /* 0x5800c00004411c03 */
/*00d0*/ ISETP.GE.U32.AND P0, PT, R2, R4, PT; /* 0x1b0e00001021dc03 */
/*00d8*/ @!P0 IADD R7, R4, R2; /* 0x480000000841e003 */
/*00e0*/ @!P0 SHL R7, R7, 0x2; /* 0x6000c0000871e003 */
/*00e8*/ @!P0 LDS R7, [R7]; /* 0xc10000000071e085 */
/*00f0*/ @!P0 IADD R5, R7, R5; /* 0x4800000014716003 */
/*00f8*/ @!P0 STS [R3], R5; /* 0xc900000000316085 */
/*0100*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0108*/ ISETP.GT.U32.AND P0, PT, R6, 0x83, PT; /* 0x1a0ec0020c61dc03 */
/*0110*/ @P0 BRA 0xc0; /* 0x4003fffea00001e7 */
/*0118*/ ISETP.GT.U32.AND P0, PT, R2, 0x1f, PT; /* 0x1a0ec0007c21dc03 */
/*0120*/ SSY 0x1f0; /* 0x6000000320000007 */
/*0128*/ @P0 NOP.S; /* 0x40000000000001f4 */
/*0130*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0138*/ LDS R4, [R3+0x80]; /* 0xc100000200311c85 */
/*0140*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*0148*/ STS [R3], R6; /* 0xc900000000319c85 */
/*0150*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0158*/ LDS R4, [R3+0x40]; /* 0xc100000100311c85 */
/*0160*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*0168*/ STS [R3], R6; /* 0xc900000000319c85 */
/*0170*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0178*/ LDS R4, [R3+0x20]; /* 0xc100000080311c85 */
/*0180*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*0188*/ STS [R3], R6; /* 0xc900000000319c85 */
/*0190*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0198*/ LDS R4, [R3+0x10]; /* 0xc100000040311c85 */
/*01a0*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*01a8*/ STS [R3], R6; /* 0xc900000000319c85 */
/*01b0*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*01b8*/ LDS R4, [R3+0x8]; /* 0xc100000020311c85 */
/*01c0*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*01c8*/ STS [R3], R6; /* 0xc900000000319c85 */
/*01d0*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*01d8*/ LDS R4, [R3+0x4]; /* 0xc100000010311c85 */
/*01e0*/ IADD R4, R5, R4; /* 0x4800000010511c03 */
/*01e8*/ STS.S [R3], R4; /* 0xc900000000311c95 */
/*01f0*/ ISETP.NE.AND P0, PT, R2, RZ, PT; /* 0x1a8e0000fc21dc23 */
/*01f8*/ @P0 BRA.U 0x218; /* 0x40000000600081e7 */
/*0200*/ @!P0 ISCADD R0, R0, c[0x0][0x24], 0x2; /* 0x4000400090002043 */
/*0208*/ @!P0 LDS R2, [RZ]; /* 0xc100000003f0a085 */
/*0210*/ @!P0 ST [R0], R2; /* 0x900000000000a085 */
/*0218*/ EXIT; /* 0x8000000000001de7 */
行からわかるように/*0130*/-/*01e8*/、合計が実行されるたびに、中間結果がすぐに共有メモリにフラッシュされ、完全なスレッドの可視性が得られます。
たぶん、マークハリスからのこれらのスライドを見てください。なぜ車輪の再発明をするのか。
www.uni-graz.at/~haasegu/Lectures/GPU_CUDA/Lit/reduction.pdf?page=35
各削減ステップは、他のステップに依存しています。したがって、最後に実行されたワープが削減フェーズで32のアクティブなスレッドに等しい場合にのみ同期を除外できます。2つのワープを使用するため、並列実行は保証されないため、64スレッドが必要になる前に、同期が必要になります。