2

パフォーマンスの問題を以下に示すコードに絞り込みました。このコードは、128,000 個の 64 バイト構造 (「ルール」) の配列を取り、それらを別の配列内に分散させます。たとえば、SCATTERSIZE が 10 の場合、コードは、インデックス 0、1、2、...、127999 で連続して格納されている「小さな」配列からこれらの構造体の 128,000 個をコピー (「分散」) し、配置します。 「大きな」配列内のインデックス 0、10、20、30、...、1279990。

ここで私が理解できないことは次のとおりです。コンピューティング機能 1.3 (Tesla C1060) のデバイスでは、SCATTERSIZE が 16 の倍数になると、パフォーマンスが大幅に低下します。コンピューティング機能 2.0 (Tesla C2075) のデバイスでは、SCATTERSIZE のたびにパフォーマンスがかなり低下します。は 24 の倍数です。

私は共有メモリを使用していないので、これは共有メモリバンクのことではないと思います。そして、合体には関係ないと思います。コマンドライン プロファイラーを使用して "gputime" エントリを調べると、1.3 デバイスではランタイムが 300% 増加し、2.0 デバイスではランタイムが 40% 増加していることがわかりました。私は困惑しています。コードは次のとおりです。

#include <stdio.h>
#include <cuda.h>
#include <stdint.h>

typedef struct{
  float a[4][4];
} Rule;

#ifndef SCATTERSIZE
#define SCATTERSIZE 96
#endif

__global__ void gokernel(Rule* b, Rule* s){
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  memcpy(&b[idx * SCATTERSIZE], &s[idx], sizeof(Rule));
}


int main(void){
  int blocksPerGrid = 1000;
  int threadsPerBlock = 128;
  int numThreads = blocksPerGrid * threadsPerBlock;
  printf("blocksPerGrid = %d, SCATTERSIZE = %d\n", blocksPerGrid, SCATTERSIZE);

  Rule* small;      
  Rule* big;        

  cudaError_t err = cudaMalloc(&big, numThreads * 128 * sizeof(Rule));
  printf("Malloc big: %s\n",cudaGetErrorString(err));

  err = cudaMalloc(&small, numThreads * sizeof(Rule));
  printf("Malloc small: %s\n",cudaGetErrorString(err));

  gokernel <<< blocksPerGrid, threadsPerBlock >>> (big, small);
  err = cudaThreadSynchronize();
  printf("Kernel launch: %s\n", cudaGetErrorString(err));
}
4

1 に答える 1

1

の実装__device__ memcpyは隠蔽されているため (コンパイラに組み込まれているため)、正確な原因を特定することは困難です。1 つの予感 (これについては njuffa のおかげです) は、多くのスレッドからのアドレスが 1 つまたはいくつかの物理 DRAM パーティションに分散されるのではなく、1 つまたはいくつかの物理 DRAM パーティションにマッピングされる、パーティション キャンピングと呼ばれるものであるということです。

SM 1_2/1_3 GPU では、メモリ アクセス ストライドによってはパーティション キャンピングが非常に悪くなる可能性がありますが、これは SM_2_0 デバイスから改善されているため、効果がそれほど顕著ではない理由が説明できます。

配列にパディングを追加して問題のあるオフセットを回避することで、この影響を回避できることがよくありますが、計算によっては価値がない場合があります。

于 2012-08-29T06:05:24.803 に答える