2

次のコードを使用した単純な cuda アプリケーションがあります。

#include <stdio.h>
#include <sys/time.h>
#include <stdint.h>
__global__
void daxpy(int n, int a, int *x, int *y) {
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  y[i] = x[i];
  int j;
  for(j = 0; j < 1024*10000; ++j) {
     y[i] += j%10;
  }
}
// debug time
void calc_time(struct timeval *start, const char *msg) {
   struct timeval end;
   gettimeofday(&end, NULL);
   uint64_t us = end.tv_sec * 1000000 + end.tv_usec - (start->tv_sec * 1000000 + start->tv_usec);
   printf("%s cost us = %llu\n", msg, us);
   memcpy(start, &end, sizeof(struct timeval));
}
void do_test() {
   unsigned long n = 1536;
   int *x, *y, a, *dx, *dy;
   a = 2.0;
   x = (int*)malloc(sizeof(int)*n);
   y = (int*)malloc(sizeof(int)*n);
   for(i = 0; i < n; ++i) {
      x[i] = i;
   }

   cudaMalloc((void**)&dx, n*sizeof(int));
   cudaMalloc((void**)&dy, n*sizeof(int));
   struct timeval start;
   gettimeofday(&start, NULL);
   cudaMemcpy(dx, x, n*sizeof(int), cudaMemcpyHostToDevice);

   daxpy<<<1, 512>>>(n, a, dx, dy); // this line 
   cudaThreadSynchronize();
   cudaMemcpy(y, dy, n*sizeof(int), cudaMemcpyDeviceToHost);
   calc_time(&start, "do_test ");
   cudaFree(dx);
   cudaFree(dy);
   free(x);
   free(y);
}
int main() {
   do_test();
   return 0;
}

gpu カーネル呼び出しは次のdaxpy<<<1, 512>>>(n, a, dx, dy)とおりで、さまざまなブロック サイズを使用していくつかのテストを実行しました。

  • daxpy<<<1, 32>>>(n, a, dx, dy)
  • daxpy<<<1, 64>>>(n, a, dx, dy)
  • daxpy<<<1, 128>>>(n, a, dx, dy)
  • daxpy<<<1, 129>>>(n, a, dx, dy)
  • daxpy<<<1, 512>>>(n, a, dx, dy)

...そして、次の観察を行いました。

  • 32、64、および128ブロックサイズの実行時間は同じです。
  • 128実行時間は、ブロック サイズとによって異なります129。特に:
    • 128実行時間は280msなので、
    • 129実行時間は 386ms です。

128ブロック サイズとの実行時間の違いの原因をお聞きしたいと思い129ます。

私の GPU は tesla K80 です:

CUDA Driver Version / Runtime Version          6.5 / 6.5
CUDA Capability Major/Minor version number:    3.7
Total amount of global memory:                 11520 MBytes (12079136768 bytes)
(13) Multiprocessors, (192) CUDA Cores/MP:     2496 CUDA Cores
GPU Clock rate:                                824 MHz (0.82 GHz)
Memory Clock rate:                             2505 Mhz
Memory Bus Width:                              384-bit
L2 Cache Size:                                 1572864 bytes
Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
Total amount of constant memory:               65536 bytes
Total amount of shared memory per block:       49152 bytes
Total number of registers available per block: 65536
Warp size:                                     32
Maximum number of threads per multiprocessor:  2048
Maximum number of threads per block:           1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
Maximum memory pitch:                          2147483647 bytes
Texture alignment:                             512 bytes
Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
Run time limit on kernels:                     No
Integrated GPU sharing Host Memory:            No
Support host page-locked memory mapping:       Yes
Alignment requirement for Surfaces:            Yes
Device has ECC support:                        Enabled
Device supports Unified Addressing (UVA):      Yes
Device PCI Bus ID / PCI location ID:           135 / 0
4

1 に答える 1

3

コメントの1つで正確な時差を提供した後、つまり：

  • 最大 128 スレッドで 280 ミリ秒、
  • 129 以上のスレッドで 386 ミリ秒、

ワープスケジューリングに関連する問題の私の理論を間接的にサポートしていると思います. K80で使用されるチップであるGK210 ホワイトペーパーを参照してください。

  • K80 SMX はクワッド ワープ スケジューラを備えています。「クワッド ワープ スケジューラ」のセクションを参照してください。
  • これは、K80 SMX が一度に最大 128 のスレッド (4 ワープ == 128 スレッド) をスケジュールできることを意味し、これらは同時に実行されます。

したがって、129 スレッドの場合、SMX は 5 つのワープをスケジュールする必要があるため、スケジュールを一度に行うことはできません。つまり、スケジュールは 2 段階で行われます。

上記が当てはまる場合、私は次のことを期待します。

  • ブロックサイズが1～128の場合、実行時間はほぼ同じです。
  • ブロックサイズが 129 ～ 192 の場合、実行時間はほぼ同じになります。

192 は SMX のコア数です。ホワイトペーパーを参照してください。注意点として、ブロック全体は常に 1 つの SMX に対してスケジュールされるため、明らかに 192 を超えるスレッドを生成する場合、それらは確実に並列実行できず、実行時間は 193 を超える数のスレッドに対して長くなるはずです。

カーネルコードをほとんど何もしない程度に単純化することで、上記の命題を検証できます。そのため、スケジューリングだけが原因で実行に時間がかかるかどうかは多かれ少なかれ明白になるはずです (メモリスループットなどの他の制限要因はありません)。 .

免責事項: K80 やクアッド ワープ スケジューラを備えた他の GPU にアクセスできないため、コードを適切にプロファイリングできないため、上記は単なる私の仮定です。とにかく、それはあなたの仕事だと思います.nvprofを使用してコードを自分でプロファイリングしてみませんか? 次に、時差がどこにあるのかを確認できるはずです。

于 2015-05-07T07:55:35.930 に答える