ひびろぐ ver.3.1

先日えいやっと最新の野良グラフィックスドライバを入れたTypeZだが、気がついたらHDMI（から変換したDVI）から映像出力ができない……いや、本体ディスプレイとのデュアルディスプレイができないことに気がついた。D-SubならOK。

一度WindowsUpdateだかVaioUpdateだかをやっちゃってドライバが飛んだのを復元などしてあるのが原因か？などとも考えたが、とにかくドライバを削除して入れ直してみた。

ドライバの削除と再インストールとシステムファイルの認証とを小一時間繰り返した結果、動作実績のある「最新ではない野良グラフィックスドライバ」をいれたらDVIデュアルディスプレイ環境が復活した。要するに、つい先日のドライバ更新を無かったことにしたわけで。

あぁ……もうグラフィックスドライバで苦労するようなノートは使いたくないな。むしろノートPC買い換えたいな。でもちょうど良いノートPCってのがどの辺なのか良くわからないな。難しいなぁ。

昨日はPage-LockedHostMemoryでカオスったわけだが、今日はオーバーラップでカオスることにする。

まぁ昨日のよりは納得がいくものが見えてる。

プログラミングガイドを読む

• 3.2.6.2 Overlap of Data Transfer and Kernel Execution
  • 「page-lockedメモリとデバイスメモリ間のコピー」と「GPUカーネル実行」は平行実行できるよ！
  • deviceOverlapプロパティが有効であることが必要だよ！
  • 「CUDA arrayやcudaMallocPitchを用いたCUDA 2D array」以外で使えるよ！
    • that do not involveってのがどこまでかかるのか（or以降までかかるのか）ちょっと不安だけど……

実験してみる

シンプルに、行列積を繰り返し実行するプログラムを書いた。

今回はカーネルがどうでも良いので、呼び出し側だけ晒す。

（転送と演算のバランスがおいしい必要はあるけど。）

シンプルな繰り返し計算

まずは単純に10回繰り返し。

  for(i=0; i<LOOP; i++){
    h_out[i] = (float*)malloc(sizeof(float)*nSize*nSize);
    h_in1[i] = (float*)malloc(sizeof(float)*nSize*nSize);
    h_in2[i] = (float*)malloc(sizeof(float)*nSize*nSize);
    cudaMalloc((void**)&d_out[i], sizeof(float)*nSize*nSize);
    cudaMalloc((void**)&d_in1[i], sizeof(float)*nSize*nSize);
    cudaMalloc((void**)&d_in2[i], sizeof(float)*nSize*nSize);
    for(int j=0; j<nSize*nSize; j++){
      h_in1[i][j] = frand();
      h_in2[i][j] = frand();
    }
  }

  EventCreate(&alls);
  EventStart(&alls);

  for(i=0; i<LOOP; i++){
#ifdef _INNER_TIME
    EventCreate(&send);
    EventCreate(&calc);
    EventCreate(&recv);
    EventStart(&send);
#endif
    cudaMemcpy(d_in1[i], h_in1[i], sizeof(float)*nSize*nSize, cudaMemcpyHostToDevice);
    cudaMemcpy(d_in2[i], h_in2[i], sizeof(float)*nSize*nSize, cudaMemcpyHostToDevice);
#ifdef _INNER_TIME
    EventStop(&send);
    EventStart(&calc);
#endif
    mmkernel<<<grid, block>>>(d_out[i], d_in1[i], d_in2[i], nSize);
    cutilSafeCall(cudaThreadSynchronize());
#ifdef _INNER_TIME
    EventStop(&calc);
    EventStart(&recv);
#endif
    cudaMemcpy(h_out[i], d_out[i], sizeof(float)*nSize*nSize, cudaMemcpyDeviceToHost);
#ifdef _INNER_TIME
    EventStop(&recv);
    EventDelete(&send);
    EventDelete(&calc);
    EventDelete(&recv);
    float fSend, fCalc, fRecv;
    fSend = GetEventTime(&send);
    fCalc = GetEventTime(&calc);
    fRecv = GetEventTime(&recv);
    printf(" send : %f nsec, %f MByte/sec\n", fSend*1000.0f, (double)(sizeof(float)*nSize*nSize)/1024.0/1024.0/(fSend/1000.0)*2.0);
    printf(" calc : %f nsec\n", fCalc*1000.0f);
    printf(" recv : %f nsec, %f MByte/sec\n", fRecv*1000.0f, (double)(sizeof(float)*nSize*nSize)/1024.0/1024.0/(fRecv/1000.0));
#endif
  }

  EventStop(&alls);
  EventDelete(&alls);

LOOPは定数。#ifdefで内部の実行時間確認をON/OFFできるようにしてある。

実行結果は以下の通り。

全体のみ測定：
 all : 22197.503906 nsec
内部も測定：
 send : 729.600037 nsec, 2741.228024 MByte/sec
 calc : 886.432007 nsec
 recv : 685.311951 nsec, 1459.189448 MByte/sec
 send : 716.607971 nsec, 2790.926188 MByte/sec
 calc : 857.280029 nsec
 recv : 677.247986 nsec, 1476.563974 MByte/sec
 send : 703.296021 nsec, 2843.752821 MByte/sec
 calc : 857.919983 nsec
 recv : 683.104004 nsec, 1463.905981 MByte/sec
 send : 686.272034 nsec, 2914.296266 MByte/sec
 calc : 863.712036 nsec
 recv : 677.824036 nsec, 1475.309180 MByte/sec
 send : 695.936035 nsec, 2873.827379 MByte/sec
 calc : 861.024048 nsec
 recv : 676.639954 nsec, 1477.890811 MByte/sec
 send : 690.720032 nsec, 2895.529212 MByte/sec
 calc : 860.319946 nsec
 recv : 683.583984 nsec, 1462.878062 MByte/sec
 send : 685.087952 nsec, 2919.333081 MByte/sec
 calc : 863.648010 nsec
 recv : 682.847961 nsec, 1464.454804 MByte/sec
 send : 682.015991 nsec, 2932.482458 MByte/sec
 calc : 860.416016 nsec
 recv : 676.736023 nsec, 1477.681111 MByte/sec
 send : 692.607971 nsec, 2887.636301 MByte/sec
 calc : 863.615967 nsec
 recv : 675.328003 nsec, 1480.761905 MByte/sec
 send : 693.888000 nsec, 2882.309501 MByte/sec
 calc : 864.416016 nsec
 recv : 680.351990 nsec, 1469.827443 MByte/sec
 all : 22779.103516 nsec

直感的に全てのsend/calc/recvを足したらallになってる気がする。

HostAllocした場合

続いてHostAlloc版、やはり10回繰り返し。

  for(i=0; i<LOOP; i++){
    cutilSafeCall(cudaHostAlloc((void**)&h_out[i], sizeof(float)*nSize*nSize, 0));
    cutilSafeCall(cudaHostAlloc((void**)&h_in1[i], sizeof(float)*nSize*nSize, 0));
    cutilSafeCall(cudaHostAlloc((void**)&h_in2[i], sizeof(float)*nSize*nSize, 0));
    cudaMalloc((void**)&d_out[i], sizeof(float)*nSize*nSize);
    cudaMalloc((void**)&d_in1[i], sizeof(float)*nSize*nSize);
    cudaMalloc((void**)&d_in2[i], sizeof(float)*nSize*nSize);
    for(int j=0; j<nSize*nSize; j++){
      h_in1[i][j] = frand();
      h_in2[i][j] = frand();
    }
  }

  EventCreate(&alls);
  EventStart(&alls);

  for(i=0; i<LOOP; i++){
#ifdef _INNER_TIME
    EventCreate(&send);
    EventCreate(&calc);
    EventCreate(&recv);
    EventStart(&send);
#endif
    cudaMemcpy(d_in1[i], h_in1[i], sizeof(float)*nSize*nSize, cudaMemcpyHostToDevice);
    cudaMemcpy(d_in2[i], h_in2[i], sizeof(float)*nSize*nSize, cudaMemcpyHostToDevice);
#ifdef _INNER_TIME
    EventStop(&send);
    EventStart(&calc);
#endif
    mmkernel<<<grid, block>>>(d_out[i], d_in1[i], d_in2[i], nSize);
    cutilSafeCall(cudaThreadSynchronize());
#ifdef _INNER_TIME
    EventStop(&calc);
    EventStart(&recv);
#endif
    cudaMemcpy(h_out[i], d_out[i], sizeof(float)*nSize*nSize, cudaMemcpyDeviceToHost);
#ifdef _INNER_TIME
    EventStop(&recv);
    EventDelete(&send);
    EventDelete(&calc);
    EventDelete(&recv);
    float fSend, fCalc, fRecv;
    fSend = GetEventTime(&send);
    fCalc = GetEventTime(&calc);
    fRecv = GetEventTime(&recv);
    printf(" send : %f nsec, %f MByte/sec\n", fSend*1000.0f, (double)(sizeof(float)*nSize*nSize)/1024.0/1024.0/(fSend/1000.0)*2.0);
    printf(" calc : %f nsec\n", fCalc*1000.0f);
    printf(" recv : %f nsec, %f MByte/sec\n", fRecv*1000.0f, (double)(sizeof(float)*nSize*nSize)/1024.0/1024.0/(fRecv/1000.0));
#endif
  }

  EventStop(&alls);
  EventDelete(&alls);

実行結果は以下の通り。

全体：
 all : 13931.423828 nsec
詳細：
 send : 370.175995 nsec, 5402.835591 MByte/sec
 calc : 884.576050 nsec
 recv : 179.008011 nsec, 5586.342290 MByte/sec
 send : 366.944000 nsec, 5450.423172 MByte/sec
 calc : 857.536011 nsec
 recv : 179.711990 nsec, 5564.458767 MByte/sec
 send : 366.272003 nsec, 5460.422818 MByte/sec
 calc : 857.760010 nsec
 recv : 178.783997 nsec, 5593.341752 MByte/sec
 send : 365.824005 nsec, 5467.109659 MByte/sec
 calc : 856.192017 nsec
 recv : 178.975998 nsec, 5587.341340 MByte/sec
 send : 369.376007 nsec, 5414.536894 MByte/sec
 calc : 859.711975 nsec
 recv : 178.688004 nsec, 5596.346560 MByte/sec
 send : 367.040009 nsec, 5448.997261 MByte/sec
 calc : 858.911987 nsec
 recv : 179.104004 nsec, 5583.348212 MByte/sec
 send : 367.488007 nsec, 5442.354627 MByte/sec
 calc : 860.032043 nsec
 recv : 178.783997 nsec, 5593.341752 MByte/sec
 send : 366.239990 nsec, 5460.900034 MByte/sec
 calc : 858.976013 nsec
 recv : 178.847992 nsec, 5591.340183 MByte/sec
 send : 366.751984 nsec, 5453.276348 MByte/sec
 calc : 858.880005 nsec
 recv : 179.360001 nsec, 5575.379055 MByte/sec
 send : 366.048004 nsec, 5463.764193 MByte/sec
 calc : 858.015991 nsec
 recv : 179.360001 nsec, 5575.379055 MByte/sec
 all : 14384.480469 nsec

転送速度が向上している。あとは単純に足しあわせたのと同じくらいっぽい。妥当。

Mappedメモリを使ってみた場合

昨日はまった（芳しい結果が得られなかった）Mappedメモリに挑戦。

  for(i=0; i<LOOP; i++){
    cutilSafeCall(cudaHostAlloc((void**)&h_out[i], sizeof(float)*nSize*nSize, cudaHostAllocMapped));
    cutilSafeCall(cudaHostAlloc((void**)&h_in1[i], sizeof(float)*nSize*nSize, cudaHostAllocMapped));
    cutilSafeCall(cudaHostAlloc((void**)&h_in2[i], sizeof(float)*nSize*nSize, cudaHostAllocMapped));
    cutilSafeCall(cudaHostGetDevicePointer((void**)&d_out[i], (void*)h_out[i], 0));
    cutilSafeCall(cudaHostGetDevicePointer((void**)&d_in1[i], (void*)h_in1[i], 0));
    cutilSafeCall(cudaHostGetDevicePointer((void**)&d_in2[i], (void*)h_in2[i], 0));
    for(int j=0; j<nSize*nSize; j++){
      h_in1[i][j] = frand();
      h_in2[i][j] = frand();
    }
  }

  EventCreate(&alls);
  EventStart(&alls);

  for(i=0; i<LOOP; i++){
#ifdef _INNER_TIME
    EventCreate(&calc);
    EventStart(&calc);
#endif
    mmkernel<<<grid, block>>>(d_out[i], d_in1[i], d_in2[i], nSize);
    //cutilSafeCall(cudaThreadSynchronize());
#ifdef _INNER_TIME
    EventStop(&calc);
    EventDelete(&calc);
    float fCalc;
    fCalc = GetEventTime(&calc);
    printf(" calc : %f nsec\n", fCalc*1000.0f);
#endif
  }

  EventStop(&alls);
  EventDelete(&alls);

これって、この範囲が終わった時点でデータはどこにあるんだろうね。このあと結果を読み出しているんだけど、そこで初めて転送されるのかしら？

実行結果は以下の通り。

全体：
 all : 80195.867188 nsec
内部：
 calc : 7945.184082 nsec
 calc : 8100.064453 nsec
 calc : 8799.391602 nsec
 calc : 8135.999512 nsec
 calc : 8501.215820 nsec
 calc : 8317.055664 nsec
 calc : 8582.559570 nsec
 calc : 8350.496094 nsec
 calc : 8150.367676 nsec
 calc : 8263.648438 nsec
 all : 83270.179688 nsec

……あれ？？？

どう見ても実行時間が延びまくってる。まぁ合計で83msecしかかかってないんだけどさ……。

Streamを使ってみた場合

Streamを使ってオーバーラップ、というのも有りみたいなので試してみた。

  for(i=0; i<LOOP; i++){
    cutilSafeCall(cudaHostAlloc((void**)&h_out[i], sizeof(float)*nSize*nSize, 0));
    cutilSafeCall(cudaHostAlloc((void**)&h_in1[i], sizeof(float)*nSize*nSize, 0));
    cutilSafeCall(cudaHostAlloc((void**)&h_in2[i], sizeof(float)*nSize*nSize, 0));
    cudaMalloc((void**)&d_out[i], sizeof(float)*nSize*nSize);
    cudaMalloc((void**)&d_in1[i], sizeof(float)*nSize*nSize);
    cudaMalloc((void**)&d_in2[i], sizeof(float)*nSize*nSize);
    for(int j=0; j<nSize*nSize; j++){
      h_in1[i][j] = frand();
      h_in2[i][j] = frand();
    }
  }

  cudaStream_t *stream;
  stream = (cudaStream_t*)malloc(sizeof(cudaStream_t)*LOOP);

  EventCreate(&alls);
  EventStart(&alls);

  for(i=0; i<LOOP; i++){
    cudaStreamCreate(&stream[i]);
  }
  for(i=0; i<LOOP; i++){
    cudaMemcpyAsync(d_in1[i], h_in1[i], sizeof(float)*nSize*nSize, cudaMemcpyHostToDevice, stream[i]);
    cudaMemcpyAsync(d_in2[i], h_in2[i], sizeof(float)*nSize*nSize, cudaMemcpyHostToDevice, stream[i]);
  }
  for(i=0; i<LOOP; i++){
    mmkernel<<<grid, block, 0, stream[i]>>>(d_out[i], d_in1[i], d_in2[i], nSize);
  }
  for(i=0; i<LOOP; i++){
    cudaMemcpyAsync(h_out[i], d_out[i], sizeof(float)*nSize*nSize, cudaMemcpyDeviceToHost);
  }
  EventStop(&alls);
  EventDelete(&alls);

実行結果は以下の通り。

全体：
 all : 9522.400391 nsec
内部：
 ……って、どうすりゃいいんだ？Asyncの測定したら意味ないよね？

HostAllocが

• send 370
• calc 860
• recv 180

だった。理想的に考えると

send+calc*10+recv=9150

なわけだが、これはあながち外れていないなぁ。

ついでにGeForce285でも実験してみた

以下、順番にまとめて貼り付け。

デフォ：
 all : 32609.570312 nsec

 send : 742.496033 nsec, 2693.617155 MByte/sec
 calc : 1901.280029 nsec
 recv : 705.120056 nsec, 1418.198266 MByte/sec
 send : 727.263977 nsec, 2750.033049 MByte/sec
 calc : 1865.216064 nsec
 recv : 698.848022 nsec, 1430.926306 MByte/sec
 send : 720.671997 nsec, 2775.187559 MByte/sec
 calc : 1869.439941 nsec
 recv : 695.743958 nsec, 1437.310320 MByte/sec
 send : 702.304016 nsec, 2847.769604 MByte/sec
 calc : 1858.239990 nsec
 recv : 688.864014 nsec, 1451.665366 MByte/sec
 send : 707.647949 nsec, 2826.263988 MByte/sec
 calc : 1865.952026 nsec
 recv : 688.480042 nsec, 1452.474976 MByte/sec
 send : 710.304016 nsec, 2815.695728 MByte/sec
 calc : 1867.167969 nsec
 recv : 696.063965 nsec, 1436.649513 MByte/sec
 send : 696.127991 nsec, 2873.034800 MByte/sec
 calc : 1860.735962 nsec
 recv : 689.408020 nsec, 1450.519857 MByte/sec
 send : 708.031982 nsec, 2824.731038 MByte/sec
 calc : 1878.559937 nsec
 recv : 686.496033 nsec, 1456.672685 MByte/sec
 send : 702.911987 nsec, 2845.306492 MByte/sec
 calc : 1866.719971 nsec
 recv : 696.447998 nsec, 1435.857321 MByte/sec
 send : 693.632019 nsec, 2883.373288 MByte/sec
 calc : 1864.255981 nsec
 recv : 692.063965 nsec, 1444.953095 MByte/sec
 all : 33175.871094 nsec

HostAlloc:
 all : 24316.193359 nsec

 send : 382.048004 nsec, 5234.944152 MByte/sec
 calc : 1881.760010 nsec
 recv : 203.743988 nsec, 4908.120133 MByte/sec
 send : 377.344025 nsec, 5300.203356 MByte/sec
 calc : 1864.031982 nsec
 recv : 203.743988 nsec, 4908.120133 MByte/sec
 send : 379.743988 nsec, 5266.706085 MByte/sec
 calc : 1858.080078 nsec
 recv : 203.328003 nsec, 4918.161824 MByte/sec
 send : 377.951996 nsec, 5291.677118 MByte/sec
 calc : 1857.471924 nsec
 recv : 205.056000 nsec, 4876.716681 MByte/sec
 send : 379.135986 nsec, 5275.151979 MByte/sec
 calc : 1867.104004 nsec
 recv : 204.255997 nsec, 4895.817059 MByte/sec
 send : 378.527985 nsec, 5283.625004 MByte/sec
 calc : 1869.567993 nsec
 recv : 204.288010 nsec, 4895.049986 MByte/sec
 send : 377.855988 nsec, 5293.021871 MByte/sec
 calc : 1865.216064 nsec
 recv : 203.871994 nsec, 4905.038571 MByte/sec
 send : 379.103973 nsec, 5275.597359 MByte/sec
 calc : 1874.752075 nsec
 recv : 203.520004 nsec, 4913.522013 MByte/sec
 send : 377.183990 nsec, 5302.451803 MByte/sec
 calc : 1866.271973 nsec
 recv : 204.160004 nsec, 4898.119006 MByte/sec
 send : 378.688019 nsec, 5281.392492 MByte/sec
 calc : 1869.119995 nsec
 recv : 204.255997 nsec, 4895.817059 MByte/sec
 all : 24895.201172 nsec

Mapped:
 all : 228410.281250 nsec

 calc : 22700.255859 nsec
 calc : 22870.496094 nsec
 calc : 24236.798828 nsec
 calc : 22738.974609 nsec
 calc : 22684.255859 nsec
 calc : 23392.416016 nsec
 calc : 22410.273438 nsec
 calc : 24593.183594 nsec
 calc : 23837.728516 nsec
 calc : 23719.263672 nsec
 all : 233330.015625 nsec

Overlap:
 all : 21263.423828 nsec
 all : 21241.408203 nsec

まぁ同じような傾向に見えるね。

つーわけで、とりあえずStreamを使ったオーバーラップはできた。Streamごとに使うデータが分かれていないと行けないはずなので、使いどころは注意が必要なんだろうな。

あと、Mappedメモリがやっぱり駄目だ。何か設定間違えるのかなあ？

cudaSetDeviceFlags(cudaDeviceMapHost);

はちゃんとやってるんだけどなあ……。

もっとでかい問題で試さないと駄目だったりするのだろうか。難しいのう。

そういえばHostAllocのみ（WriteCombinedなどの設定をしない）でのテストはやってなかった。

./ha 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy           : 16029.792786 nsec, 5864.080794 MByte/sec
  all              : 1158009.052277 nsec
part2(indirect)
  memcpy           : 16031.936646 nsec, 5863.296623 MByte/sec
  all              : 78474.044800 nsec
part3(indirect random)
  memcpy           : 16034.784317 nsec, 5862.255341 MByte/sec
  all              : 78482.151031 nsec
./ha_wk 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy+kernel    : 58448.158264 nsec, 1608.262823 MByte/sec
  all              : 1196021.080017 nsec
part1(CPU)
  CPUkernel        : 14250.993729 nsec
  memcpy+CPUkernel : 30297.216415 nsec, 3102.595259 MByte/sec
  all              : 84782.838821 nsec
part2(indirect)
  memcpy+kernel    : 96689.697266 nsec, 972.182173 MByte/sec
  all              : 160046.100616 nsec
part2(CPU)
  CPUkernel        : 20169.973373 nsec
  memcpy+CPUkernel : 36217.407227 nsec, 2595.437034 MByte/sec
  all              : 98576.068878 nsec
part3(indirect random)
  memcpy+kernel    : 165139.709473 nsec, 569.215002 MByte/sec
  all              : 227524.995804 nsec
part3(CPU)
  CPUkernel        : 153204.917908 nsec
  memcpy+CPUkernel : 169245.315552 nsec, 555.406805 MByte/sec
  all              : 231613.159180 nsec

• 転送速度は十分に得られた
• CPUkernelの実行時間が短くなったりならなかったりのちょっと謎の挙動。なんだこれ？

プログラム内部で一発しか実行していないけど、何度か実験して同じような値が見える。うーん、HostAlloc周りも対象プログラムによって色々調整しないと行けないのか……これ、チューニング無理くせー。

次はオーバーラップの実験かな。

なんだか久々にCUDAネタ。twitterでぶつぶつ言いながらやった実験なので、こちらに投稿しておく。（後であちらにも投げるつもりだけど。）

メチャクチャ長くなったけど仕方がない。

Well-KnownなCUDA高速化技術の１つにPage-Lockedなメモリを使うっていうのがある。というわけで、PinnedMemoryに関する実験をやってみた。

一応真面目にやってはいるんだけど、まだ理解が追いついていない部分もあるので、参考にする際は自己責任で。むしろ情報交換しましょう。間違いへの突っ込みも歓迎。

実験に使ったプログラムのフルソースコードなどは末尾。

とりあえずProgrammingGuideを読んでみる

3.0のProgrammingGuideだと3.2.5項。

以下要約：

• 3.2.5 Page-Locked Host Memory
  • OSのメモリ制御をコントロールすることで性能を上げることができるよ！（制御をコントロールって酷い言葉だな！）
  • page-lockedメモリはpinnedメモリとも呼ばれるよ！
  • host-device間のデータコピーとカーネル実行が平行にできるようになるよ！
  • host-device間で同じアドレスが使えるようになるよ！
  • host-device間のデータ転送速度を向上させられるよ！
• 確保しすぎるとOSの制御を阻害するとか、副作用もあるので気をつけようね！
• 3.2.5.1 Portable Memory
  • 複数スレッド（OS側のスレッド）で使うときはcudaHostAllocPortableを設定すると良いよ！
• 3.2.5.2 Write-Combining Memory
  • cudaHostAllocWriteCombinedを設定するとOSのキャッシュ管理から切り離されるよ！
  • host-device間の転送速度が上がるよ！
  • hostでWCメモリを読むと遅いよ！だからhostからは書き込むだけにした方が良いよ！（device側の読み書き速度がどうなるのかは書かれてないなあ）
• 3.2.5.3 Mapped Memory
  • cudaHostAllocMappedを設定すると、hostとdeviceで同じメモリ空間を使えるよ！
  • host-device間でメモリコピーを明示的に書かなくても良くなるし、計算と転送のオーバーラップも簡単だよ！
  • ただし、read-after-write/write-after-read/write-after-writeの問題があるのでstreamsかeventsでタイミング制御をする必要はあるよ！
  • atomic演算は注意してね！

だいたいこんなかんじ。間違ってたらごめん。

考えてみる

なんとなくわかった。

• WriteCombinedにすれば転送速度が上がる。

OK、OK。ただしOSのキャッシュ制御の都合もあるから、CPUでもGPUでもいじりまくるようなデータは気をつけろということだな。

• Mappedにすればアドレスの一元化が可能になる。memcpyとか不要。通信と演算のオーバーラップとか楽勝。

これは便利かもしれない。でも、それってつまり必要に応じてデータがコピーされるってことだよね？アクセスパターンとか次第で酷い性能になるよね？？？

テストしてみる

まずはサンプルに入っているbandwidthTestを走らせる。

>/usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest --memory=pageable
[bandwidthTest]
/usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest Starting...

Running on...

 Device 0: GeForce GTX 480
 Quick Mode

 Host to Device Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     5395.7

 Device to Host Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     4587.9

 Device to Device Bandwidth, 1 Device(s)
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     111025.4

[bandwidthTest] - Test results:
PASSED

Press  to Quit...
-----------------------------------------------------------

>/usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest --memory=pinned
[bandwidthTest]
/usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest Starting...

Running on...

 Device 0: GeForce GTX 480
 Quick Mode

 Host to Device Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     5815.4

 Device to Host Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     6105.8

 Device to Device Bandwidth, 1 Device(s)
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     110967.5

[bandwidthTest] - Test results:
PASSED

Press  to Quit...
-----------------------------------------------------------

H2Dで10%、D2Hで30%くらい高速化されてる。これはすごい。

ソースを眺めたところ、どうやら本当にcudaHostAllocにcudaHostAllocWriteCombinedを指定しているだけっぽい。

コードを書いてテストしてみる

以下の比較をしてみよう：

• 普通にMalloc/cudaMalloc
• cudaHostAlloc with cudaHostAllocWriteCombined
• HostAllocMapped

それぞれについて、

• 単純に配列を捜査するカーネル
• 間接アクセスがあるカーネル
• 間接アクセスでしかも乱数なカーネル（WriteCombinedが必要に応じてデータをやりとりする機能の場合にうぼあしそうなカーネル）

を比較。

ついでに、カーネルをCPUにやらせた版も作成して様子見。

GPUの計算は、いずれもgridDim=(1,1,1) blockDim=(128,1,1)という適当な設定。これでも十分に傾向くらいは掴めると思いたい。

カーネルは以下の通り：

// simpleテスト用(part1)
__global__ void kernel1
(float *data1, float *data2, int size)
{
  int i;
  int begin, end, step;
  begin = threadIdx.x;
  end = size;
  step = blockDim.x;
  for(i=begin; i
カーネル以外は添付ファイルにて。
基本的にカーネルを一発実行するだけのプログラム。

まずは普通にMalloc/cudaMallocした場合：
./normal 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy           : 34878.208160 nsec, 2695.092580 MByte/sec
  all              : 1128581.047058 nsec
part2(indirect)
  memcpy           : 34505.214691 nsec, 2724.225913 MByte/sec
  all              : 49117.088318 nsec
part3(indirect random)
  memcpy           : 34498.271942 nsec, 2724.774161 MByte/sec
  all              : 49160.003662 nsec
./normal_wk 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy+kernel    : 77812.156677 nsec, 1208.037459 MByte/sec
  all              : 1167029.857635 nsec
part1(CPU)
  CPUkernel        : 24100.780487 nsec
  memcpy+CPUkernel : 43626.399994 nsec, 2154.658647 MByte/sec
  all              : 48449.993134 nsec
part2(indirect)
  memcpy+kernel    : 115774.589539 nsec, 811.922550 MByte/sec
  all              : 130471.944809 nsec
part2(CPU)
  CPUkernel        : 28672.933578 nsec
  memcpy+CPUkernel : 48061.695099 nsec, 1955.819490 MByte/sec
  all              : 62689.065933 nsec
part3(indirect random)
  memcpy+kernel    : 184142.211914 nsec, 510.475024 MByte/sec
  all              : 198773.860931 nsec
part3(CPU)
  CPUkernel        : 126466.035843 nsec
  memcpy+CPUkernel : 145819.137573 nsec, 644.634179 MByte/sec
  all              : 160428.047180 nsec

_wkがカーネルも実行している版。流石に間接参照やランダム参照が遅い。でも実はもっと遅くなるんじゃないかと思っていた。
これが基準値ということで。
（CPUがkernelを実行した場合にMByte/secを表示する意味は無さそうだなぁ。）


WriteCombined設定をした場合：
./wc 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy           : 16025.951385 nsec, 5865.486406 MByte/sec
  all              : 1158394.098282 nsec
part2(indirect)
  memcpy           : 16022.495270 nsec, 5866.751615 MByte/sec
  all              : 82098.960876 nsec
part3(indirect random)
  memcpy           : 16023.904800 nsec, 5866.235551 MByte/sec
  all              : 82008.123398 nsec
./wc_wk 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy+kernel    : 58642.913818 nsec, 1602.921715 MByte/sec
  all              : 1202525.138855 nsec
part1(CPU)
  CPUkernel        : 1574666.976929 nsec
  memcpy+CPUkernel : 1590638.549805 nsec, 59.095764 MByte/sec
  all              : 1648941.993713 nsec
part2(indirect)
  memcpy+kernel    : 97171.073914 nsec, 967.366071 MByte/sec
  all              : 163262.128830 nsec
part2(CPU)
  CPUkernel        : 1573044.061661 nsec
  memcpy+CPUkernel : 1589031.250000 nsec, 59.155539 MByte/sec
  all              : 1657025.098801 nsec
part3(indirect random)
  memcpy+kernel    : 165483.520508 nsec, 568.032392 MByte/sec
  all              : 231626.033783 nsec
part3(CPU)
  CPUkernel        : 1745186.090469 nsec
  memcpy+CPUkernel : 1761167.358398 nsec, 53.373690 MByte/sec
  all              : 1827242.136002 nsec

カーネル非実行のmemcpyをみてわかるように、転送速度がとても速くなった。ただしmalloc等の処理に手間取るようで、allの値はむしろ増加。
CPUkernelの値が増加。ProgrammingGuideにあったhostから叩くデータは良くない的な記述が反映されている模様。注意が必要だなあ。


つづいてMapped：
./mapped 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy           : 2.528000 nsec, 37183545.414364 MByte/sec
  all              : 1177767.992020 nsec
part2(indirect)
  memcpy           : 2.336000 nsec, 40239724.436970 MByte/sec
  all              : 101921.796799 nsec
part3(indirect random)
  memcpy           : 2.336000 nsec, 40239724.436970 MByte/sec
  all              : 101827.144623 nsec
./mapped_wk 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy+kernel    : 141869.567871 nsec, 662.580435 MByte/sec
  all              : 1318675.041199 nsec
part1(CPU)
  CPUkernel        : 17521.142960 nsec
  memcpy+CPUkernel : 17529.504776 nsec, 5362.387654 MByte/sec
  all              : 109644.174576 nsec
part2(indirect)
  memcpy+kernel    : 158290.374756 nsec, 593.845331 MByte/sec
  all              : 260251.045227 nsec
part2(CPU)
  CPUkernel        : 20175.933838 nsec
  memcpy+CPUkernel : 20186.048508 nsec, 4656.681567 MByte/sec
  all              : 123075.962067 nsec
part3(indirect random)
  memcpy+kernel    : 340836.029053 nsec, 275.792440 MByte/sec
  all              : 442532.062531 nsec
part3(CPU)
  CPUkernel        : 152827.024460 nsec
  memcpy+CPUkernel : 152834.014893 nsec, 615.046330 MByte/sec
  all              : 255613.803864 nsec

カーネル非実行：明示的なmallocが消えたのでmemcpyの値が吹っ飛んだ。allの値がさらに増えていることから、やはりオーバーヘッドは小さくないのだろうか。
カーネル実行：困ったことに、何も速くなっている気がしない。


最後に、2つのパラメタを両方設定してみた：
./mpwc 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy           : 2.528000 nsec, 37183545.414364 MByte/sec
  all              : 1180523.872375 nsec
part2(indirect)
  memcpy           : 2.336000 nsec, 40239724.436970 MByte/sec
  all              : 106031.894684 nsec
part3(indirect random)
  memcpy           : 2.336000 nsec, 40239724.436970 MByte/sec
  all              : 105930.805206 nsec
./mpwc_wk 0 50000000
size = 50000000Byte(48828KByte, 47MByte), 12500000 elements
device max=2, n=0
part1(simple)
  memcpy+kernel    : 141808.486938 nsec, 662.865827 MByte/sec
  all              : 1324396.133423 nsec
part1(CPU)
  CPUkernel        : 1578817.129135 nsec
  memcpy+CPUkernel : 1578768.432617 nsec, 59.540081 MByte/sec
  all              : 1676369.905472 nsec
part2(indirect)
  memcpy+kernel    : 158298.950195 nsec, 593.813161 MByte/sec
  all              : 264221.906662 nsec
part2(CPU)
  CPUkernel        : 1572522.163391 nsec
  memcpy+CPUkernel : 1572480.102539 nsec, 59.778181 MByte/sec
  all              : 1678436.040878 nsec
part3(indirect random)
  memcpy+kernel    : 340668.151855 nsec, 275.928347 MByte/sec
  all              : 446609.020233 nsec
part3(CPU)
  CPUkernel        : 1745092.153549 nsec
  memcpy+CPUkernel : 1745044.189453 nsec, 53.866831 MByte/sec
  all              : 1851984.977722 nsec

まぁ、memcpyが要らないプログラムにmemcpy高速化をしているようなものなので、どうにもならんわな。


結論
というわけで、とりあえず書いて走らせてはみたんだけど、今ひとつだ。

WriteCombined設定で通信速度が向上
ただしhostから叩く際の速度低下があり得るので対象データに注意
Mappedが使い物にならん（コーディングが楽なのは事実）

という感じ。
もちろん、使い方が悪い可能性は否定しないので、今後追加調査を要すると言うことで。特に、通信と演算のオーバーラップがまだ試せていない。Mappedの本懐はそちらかなあと思っている。あとはPortable設定で複数GPU実験だな！

ソースコードなど
えいやっと公開。

後輩某氏に聞かれたので、せっかくだからまとめておく。

WindowsならVisualProfilerで簡単なんだけどね。

環境はCUDA_2.3。/usr/local/cuda/doc/CUDA_Profiler_2.3.txtに書かれている内容の日本語抜粋だと思ってくれればOK。

必要な知識

環境変数の設定方法。

bashならexport HOGE=FUGA、tcshならsetenv HOGE FUGA。詳しくはググれ。

実際の使い方

1. 環境変数CUDA_PROFILEに1を設定する。
2. 必要に応じて環境変数CUDA_PROFILE_LOGにプロファイル出力先を設定する。設定しなければカレントディレクトリの./cuda_profile.log.に出力される。複数GPUの場合は%dとかで制御できる模様。
3. 必要に応じて環境変数CUDA_PROFILE_CSVを設定する。1にするか０にするかで出力ファイルの書式が変わる。取り込んで使うときはお好みのものを使いましょう。
4. 環境変数CUDA_PROFILE_CONFIGに設定ファイルのパスを指定する。この設定ファイルに書かれている項目のプロファイル情報が、前述の出力先ファイルに吐かれることになる。

CUDA_PROFILE_CONFIGで指定した設定ファイルには、プロファイル項目を改行区切りで並べておけば良い。ただし、最大4つしか指定できない。（後者のカーネル実行カウンタが4つで、前者のオプションは全部OKかも？）

• オプション
  • timestamp : カーネルの実行やデータ転送の時刻。（単位はなんだろう？）
  • gridsize : GridあたりのBlock数
  • threadblocksize : BlockあたりのThread数
  • dynsmemperblock : 動的に確保された、BlockあたりのSharedMemoryの大きさ
  • stasmemperblock : 静的に確保された、BlockあたりのSharedMemoryの大きさ
  • regperthread : Threadあたりのレジスタ使用数
  • memtransferdir : メモリ転送の方向、0だとホストからGPU、1だとGPUからCPU
  • memtransfersize : メモリコピーのバイト数
  • streamid : ストリームのID
• カーネル実行カウンタ
  • gld_incoherent : コアレスドでないGlobalMemory読み込みの回数
  • gld_coherent : コアレスドなGlobalMemory読み込みの回数
  • gld_32b : 32byte単位のGlobalMemory読み込み回数
  • gld_64b : 64byte単位のGlobalMemory読み込み回数
  • gld_128b : 128byte単位のGlobalMemory読み込み回数
  • gld_request : GlobalMemory読み込み回数
  • gst_incoherent : コアレスドでないGlobalMemory書き込み回数
  • gst_coherent : コアレスドなGlobalMemory書き込み回数
  • gst_32b : 32byte単位のGlobalMemory書き込み回数
  • gst_64b : 64byte単位のGlobalMemory書き込み回数
  • gst_128b : 128byte単位のGlobalMemory書き込み回数
  • gst_request : GlobalMemory書き込み回数–local_load : LocalMemory読み込み回数
  • local_store : LocalMemory書き込み回数
  • branch : Threadあたりの分岐回数
  • divergent_branch : Threadあたりのdivergentな分岐回数
  • instructions : 実行された命令数
  • warp_serialize : SharedMemoryやConstantMemoryのアドレスコンフリクトによりwarpがシリアライズされた回数
  • cta_launched : Threadが実行された回数

ちょっと翻訳を勘違いしているのとかあったらごめんなさい。あと、めんどくさくて実行確認してませんごめんなさい（酷い