Ron Farberのコードの動作を理解するのに助けが必要です：http ：//www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/208801731？pgno = 2
共有メモリを使用すると、非共有メモリバージョンよりもパフォーマンスが向上することを理解していません。つまり、インデックス計算ステップをさらにいくつか追加し、別のRd / Wrサイクルを追加して共有メモリにアクセスする場合、グローバルメモリのみを使用するよりもこれをどのように高速化できますか？どちらの場合も、同じ数またはRd/Wrサイクルがグローバルメモリにアクセスします。データは、カーネルインスタンスごとに1回だけアクセスできます。データは引き続きグローバルメモリを使用して出入りします。カーネルインスタンスの数は同じです。レジスタ数は同じように見えます。処理ステップを追加すると、どのように高速化できますか。（プロセスステップを差し引くことはありません。）基本的に、より多くの作業を行っており、より速く完了しています。
共有メモリアクセスはグローバルよりもはるかに高速ですが、ゼロ（または負）ではありません。私は何が欠けていますか？
「スロー」コード：
__global__ void reverseArrayBlock(int *d_out, int *d_in) {
int inOffset = blockDim.x * blockIdx.x;
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int in = inOffset + threadIdx.x;
int out = outOffset + (blockDim.x - 1 - threadIdx.x);
d_out[out] = d_in[in];
}
「高速」コード：
__global__ void reverseArrayBlock(int *d_out, int *d_in) {
extern __shared__ int s_data[];
int inOffset = blockDim.x * blockIdx.x;
int in = inOffset + threadIdx.x;
// Load one element per thread from device memory and store it
// *in reversed order* into temporary shared memory
s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];
// Block until all threads in the block have written their data to shared mem
__syncthreads();
// write the data from shared memory in forward order,
// but to the reversed block offset as before
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int out = outOffset + threadIdx.x;
d_out[out] = s_data[threadIdx.x];
}