Compute Capability 3.5のNon-Coherent Cacheの性能評価

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

CUDA 5の新機能(5): Texture Cacheを自動的に使うDevice Memory読み込みの記事で話題にしたNon-Coherent Cache。実際に性能にどれほど寄与するのか、簡単なテストを行ってみた。

目次

おさらい

Tesla K20シリーズに代表されるCompute Capability 3.5のGPUは、Global MemoryとしてのアクセスをTexture Cache経由で行う命令(Non-Coherent Cache命令)を持っている。CUDA Cのプログラムにおいては、const修飾子と__restrict__修飾子とを付けて宣言されたGlobal Memoryへのポインタへのアクセスを記述すると、Non-Coherent Cache命令を使うような中間コードへとコンパイルされる。

テストプログラム

#include <cstdio>
#include <cstdlib>

#define KERNEL(F, R)                                            \
  template<typename MAP_T, typename IN_T, typename OUT_T>       \
  static __global__ void                                        \
  F(const MAP_T *R map,                                         \
    const IN_T *R in,                                           \
    OUT_T *out,                                                 \
    size_t n)                                                   \
  {                                                             \
    size_t i = blockIdx.x * blockDim.x + threadIdx.x;           \
    if (i >= n) {                                               \
      return;                                                   \
    }                                                           \
    out[i] = in[map[i]];                                        \
  }                                                             \

KERNEL(kernel_norestrict, )
KERNEL(kernel_restrict, __restrict)

static texture<int> t_map, t_in;

template<typename OUT_T>
static __global__ void
kernel_texture(OUT_T *out,
               size_t n)
{
  size_t i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i >= n) {
    return;
  }
  out[i] = tex1Dfetch(t_in, tex1Dfetch(t_map, i));
}

int
main(int argc, char **argv)
{
  size_t n = std::atoll(argv[1]);
  size_t niter = std::atoll(argv[2]);
  int *h_map = new int[n];
  int *h_in = new int[n];
  std::srand(0);
  for (size_t i = 0; i < n; ++i) {
    h_map[i] = std::rand() % n;
    h_in[i] = std::rand();
  }
  int *d_map, *d_in, *d_out;
  cudaMalloc(reinterpret_cast<void **>(&d_map), n * sizeof(int));
  cudaMemcpy(d_map, h_map, n * sizeof(int), cudaMemcpyHostToDevice);
  cudaMalloc(reinterpret_cast<void **>(&d_in), n * sizeof(int));
  cudaMemcpy(d_in, h_in, n * sizeof(int), cudaMemcpyHostToDevice);
  cudaMalloc(reinterpret_cast<void **>(&d_out), n * sizeof(int));
  int block = 256;
  int grid = (n + block - 1) / block;
  cudaBindTexture(0, t_map, d_map);
  cudaEvent_t start, end;
  cudaEventCreate(&start);
  cudaEventCreate(&end);
  {
    kernel_norestrict<<<grid, block>>>(d_map, d_in, d_out, n);
    int *d_tmp = d_in;
    d_in = d_out;
    d_out = d_tmp;
  }
  {
    kernel_restrict<<<grid, block>>>(d_map, d_in, d_out, n);
    int *d_tmp = d_in;
    d_in = d_out;
    d_out = d_tmp;
  }
  {
    cudaBindTexture(0, t_in, d_in);
    kernel_texture<<<grid, block>>>(d_out, n);
    int *d_tmp = d_in;
    d_in = d_out;
    d_out = d_tmp;
  }
  cudaEventRecord(start);
  for (size_t iter = 0; iter < niter; ++iter) {
    kernel_norestrict<<<grid, block>>>(d_map, d_in, d_out, n);
    int *d_tmp = d_in;
    d_in = d_out;
    d_out = d_tmp;
  }
  cudaEventRecord(end);
  cudaEventSynchronize(end);
  float ms_norestrict;
  cudaEventElapsedTime(&ms_norestrict, start, end);
  cudaEventRecord(start);
  for (size_t iter = 0; iter < niter; ++iter) {
    kernel_restrict<<<grid, block>>>(d_map, d_in, d_out, n);
    int *d_tmp = d_in;
    d_in = d_out;
    d_out = d_tmp;
  }
  cudaEventRecord(end);
  cudaEventSynchronize(end);
  float ms_restrict;
  cudaEventElapsedTime(&ms_restrict, start, end);
  cudaEventRecord(start);
  for (size_t iter = 0; iter < niter; ++iter) {
    cudaBindTexture(0, t_in, d_in);
    kernel_texture<<<grid, block>>>(d_out, n);
    int *d_tmp = d_in;
    d_in = d_out;
    d_out = d_tmp;
  }
  cudaEventRecord(end);
  cudaEventSynchronize(end);
  float ms_texture;
  cudaEventElapsedTime(&ms_texture, start, end);
  std::printf("%d\t%f\t%f\t%f\n", n, ms_norestrict, ms_restrict, ms_texture);
  return 0;
}

ホストの標準Cライブラリのrand関数を使って作った添字配列に対する、メモリギャザの処理である。ポインタ経由のGlobal Memoryアクセスに__restrict__を付ける/付けないの2パターンに、さらに従来のTextureを使った、計3パターンの処理を用意している。今回は、2の冪でデータサイズを増やしながら、10セットの処理を行った合計時間をテストする。

テスト環境

CPU: クアッドコア Xeon W3520 2.67 GHz
GPU: Tesla K20X
OS: CentOS 6.4
CUDA: CUDA 5.0
GPUドライバ: バージョン319.21開発用ドライバ

結果

ファイル:Ncgather.png

log(N) w/o __restrict__ w/ __restrict__ texture
0 0.051936 0.056288 0.126144
1 0.051968 0.055840 0.127232
2 0.051808 0.055584 0.119968
3 0.051808 0.054112 0.128512
4 0.052032 0.056256 0.124864
5 0.051904 0.054336 0.237152
6 0.052448 0.056352 0.127904
7 0.052992 0.056128 0.125760
8 0.054304 0.056736 0.127552
9 0.055680 0.058080 0.127328
10 0.057568 0.058432 0.127424
11 0.058720 0.058272 0.124960
12 0.067264 0.061536 0.133056
13 0.075616 0.064224 0.131936
14 0.091744 0.070080 0.128512
15 0.130944 0.085344 0.137088
16 0.200896 0.112704 0.160512
17 0.343392 0.174048 0.210656
18 0.670400 0.353440 0.391584
19 1.820608 1.457920 1.485440
20 4.878496 4.355488 4.376704
21 10.994880 10.447712 10.488736
22 23.269217 22.649920 22.769440
23 47.856510 47.295616 47.199295
24 97.196259 96.160385 96.187164
25 195.845215 194.337021 194.427231
26 472.194916 472.976166 473.025574
27 1396.915649 1399.452881 1399.394409
  • Texture使用の場合に生じている約7 nsのオーバーヘッドが、Non-Coherent Cacheの場合は生じない。
  • メモリアクセスに要する時間は、Non-Coherent CacheはTextureにわずかに劣る。
  • Non-Coherent CacheやTextureが目に見えて速いのは、32 KiBから4 MiB程度の範囲である。
  • その範囲より小さな場合は、Non-Coherent Cacheを使うと逆に遅くなるため、使わないほうが良い可能性がある。

というような結論が導き出せる。少なくとも、Textureが有用となるケースではNon-Coherent Cacheもやはり有用であり、オーバーヘッドの適性によってはTextureを置き換える使い方が最適解となるようである。


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:Compute Capability 3.5のNon-Coherent Cacheの性能評価

個人用ツール