3

1Dの場合、CUDAのグローバルメモリの合体​​したアクセス要件全体をほぼ理解しました。

ただし、2次元の場合には少し行き詰まっています（つまり、2Dブロックで構成された2Dグリッドがあります）。

ベクトルがin_vectorあり、カーネルに合体してアクセスしたいとします。そのようです：

__global__ void my_kernel(float* out_matrix, float* in_vector, int size)
{
   int i = blockIdx.x * blockDim.x + threadIdx.x;
   int j = blockIdx.y * blockDim.y + threadIdx.y;
   // ...
   float vx = in_vector[i]; // This is good. Here we have coalesced access
   float vy = in_vector[j]; // Not sure about this. All threads in my warp access the same global address. (See explanation)
   // ...
   // Do some calculations... Obtain result
}

この2Dの場合の私の理解では、ブロック内のスレッドは列を中心に「配置」されています。例：（threadIdx.x、threadIdx.y）表記を想定：

  • 最初のワープは次のようになります：（0、0）、（1、0）、（2、0）、...、（31、0）、
  • 2番目のワープは次のようになります：（0、1）、（1、1）、（2、1）、...、（31、1）、
  • 等々...

この場合in_vector[i]、同じワープ内の連続する各スレッドが連続するアドレスにアクセスするため、呼び出しによって合体したアクセスが得られます。ただしin_vector[j]、連続する各スレッドがグローバルメモリ内の同じアドレスにアクセスするため、呼び出しは悪い考えのようです（たとえば、ワープ0のすべてのスレッドはin_vector [0]にアクセスし、32の異なるグローバルメモリ要求が発生します）

私はこれを正しく理解しましたか？もしそうなら、どうすればグローバルメモリに合体してアクセスできますin_vector[j]か？

4

1 に答える 1

7

あなたがあなたの質問で示したことは、特定のブロックサイズに対してのみ正しいです。「合体」アクセス：

int i = blockIdx.x * blockDim.x + threadIdx.x;
float vx = in_vector[i];

が32以上の場合in_vectorにのみ、グローバルメモリからの合体アクセスが発生します。合体した場合でも、同じ値を共有するブロック内の各スレッドは、グローバルメモリから同じ単語を読み取ります。これは直感に反しているようです。そして無駄です。blockDim.xthreadIdx.x

読み取りがスレッドごとに一意で合体するようにする正しい方法は、ブロック内のスレッド番号とグリッド内のオフセットを計算することです。おそらく次のようになります。

int tid = threadIdx.x + blockDim.x * threadIdx.y; // must use column major order
int bid = blockIdx.x + gridDim.x * blockDim.y; // can either use column or row major
int offset = (blockDim.x * blockDim.y) * bid; // block id * threads per block
float vx = in_vector[tid + offset];

スレッドごとに一意の値を読み取ることを意図していない場合は、次のように、メモリ帯域幅を大幅に節約し、共有メモリを使用して合体を実現できます。

__shared__ float vx[32], vy[32]; 

int tid = threadIdx.x + blockDim.x * threadIdx.y;

if (tid < 32) {
    vx[tid] = in_vector[blockIdx.x * blockDim.x + tid];
    vy[tid] = in_vector[blockIdx.y * blockDim.y + tid];
}
__syncthread();

一意の値を共有メモリに一度読み取る単一のワープを取得します。その後、他のスレッドは、グローバルメモリアクセスをさらに必要とせずに、共有メモリから値を読み取ることができます。上記の例では、in_vectorそのように2回読むことが必ずしもそれほど意味がない場合でも、コードの規則に従っていることに注意してください。

于 2012-09-09T12:25:33.603 に答える