CUDA 通信と計算のオーバーラップを試した

マルチポスト元→http://exth.net/~tgbt/wordpress/2010/05/19/3155/

昨日は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);
はちゃんとやってるんだけどなあ……。
もっとでかい問題で試さないと駄目だったりするのだろうか。難しいのう。