プログラミング > CUDA > メモ


※上記の広告は60日以上更新のないWIKIに表示されています。更新することで広告が下部へ移動します。

CUDAメモ



メモ

nvccの-archと-code(CUDA COMPILER DRIVER NVCC v7.5より)

  • --gpu-architecture arch (短い名前-arch)
    • CUDA入力ファイルがコンパイルされなくてはならないNVIDIAの「仮想」GPUアーキテクチャのクラス名を指定する。
    • 以下の略記法として説明される例外を除き、このオプションで指定されたアーキテクチャは、「仮想」アーキテクチャでなければならない(compute_20のように)。通常、このオプション単体は、「実際の」アーキテクチャのために生成されたPTXのアセンブリには影響しない(それは、nvccのオプション--gpu-codeの役割である。以下を参照)。それよりも、その目的は、前処理と入力からPTXのコンパイルの制御することである。
    • 利便性のため、単純なnvccコンパイルの場合、以下の略記法がサポートされる。もし、オプション--gpu-codeの値が指定されなければ、そのとき、このオプションの--gpu-architectureの値を初期値とする。この状況では、上の記載の唯一の例外として、--gpu-architectureに対する指定された値が、「実際の」アーキテクチャでもよい(sm_20のような)。この場合ではnvccは指定された「実際の」アーキテクチャとそれの最も近い「仮想」アーキテクチャを有効なアーキテクチャの値として使う。例えば、nvcc --gpu-architectue=sm_20は、nvcc --gpu-architecture=compute_20 --gpu-code=sm_20,compute_20と等価である。
    • サポートされた「仮想」アーキテクチャのリストは、Virtual Architecture Feature Listを参照せよ。サポートされた「実際の」アーキテクチャのリストは、 GPU Feature Listを参照せよ。
  • --gpu-code code,... (短い名前-code)
    • アセンブルとPTXを最適化するためのNVIDIA GPU名を指定する。
    • nvccは、それぞれの指定されたコードのアーキテクチャに対して、コンパイルされたコードイメージを結果の実行ファイルに組み込み、それは、それぞれの「実際の」アーキテクチャ(sm_20のような)に対する真のバイナリロードイメージと「仮想」アーキテクチャ(compute_20のような)PTXコードである。
    • 実行中、そのような組み込まれたPTXコードは、もし現在のGPUに対するバイナリロードイメージがない場合、CUDAランタイムシステムにより動的にコンパイルされる。
    • オプション--gpu-architectureと--gpu-codeで指定されたアーキテクチャは、「仮想」であり「実際の」でもあるが、codeのアーキテクチャはarchのアーキテクチャと互換性がなければならない。--gpu-codeオプションが使われるとき、--gpu-architectureオプションは、「仮想」PTXアーキテクチャでなければならない。
    • 例えば、--gpu-architecture=compute_35は--gpu-code=sm_30と互換性がない、なぜなら先のコンパイルステージがsm_30では提供されないcompute_35の特徴の利用可能を仮定するためである。
    • サポートされた「仮想」アーキテクチャのリストは、Virtual Architecture Feature Listを参照せよ。サポートされた「実際の」アーキテクチャのリストは、 GPU Feature Listを参照せよ。

デバイスからデバイスへのmemcpyの注意

cudaMemcpy(dst, src, count, cudaMemcpyDeviceToDevice)は、ブロッキングではない。そのため必要に応じて、ホスト側でcudaThreadSynchronize()を呼ぶ。

CUDAでdouble演算

CUDAで倍精度浮動小数点数を使った場合に、どんな実行コードが生成されるのか確認する。 まずは確認のためのテストプログラム。

  • test.cu
  1. #include <stdio.h>
  2.  
  3. __global__ void mulArray(double* inputOnGPU, double* outputOnGPU)
  4. {
  5. int i = blockDim.x * blockIdx.x + threadIdx.x;
  6.  
  7. outputOnGPU[i] = 5.0 * inputOnGPU[i];
  8. }
  9.  
  10. #define BLOCK_DIM 8
  11. #define THREAD_DIM 32
  12. #define N (BLOCK_DIM * THREAD_DIM)
  13.  
  14. int main()
  15. {
  16. double input[N], output[N];
  17. double* inputOnGPU;
  18. double* outputOnGPU;
  19.  
  20. cudaMalloc((void**)&inputOnGPU, sizeof(double) * N);
  21. cudaMalloc((void**)&outputOnGPU, sizeof(double) * N);
  22.  
  23. for (int i = 0; i < N; ++i) {
  24. input[i] = (double)i;
  25. }
  26.  
  27. cudaMemcpy(inputOnGPU, input, sizeof(double) * N, cudaMemcpyHostToDevice);
  28.  
  29. mulArray<<<BLOCK_DIM, THREAD_DIM>>>(inputOnGPU, outputOnGPU);
  30.  
  31. cudaMemcpy(output, outputOnGPU, sizeof(double) * N, cudaMemcpyDeviceToHost);
  32.  
  33. for (int i = 0; i < N; ++i) {
  34. printf("%d, %lf\n", i, output[i]);
  35. }
  36. }

入力を5倍する単純なプログラム。まずは中間言語PTXを確認する。コマンドで以下を入力(環境は CUDA 2.3で32ビット)。

nvcc -arch sm_13 -O3 -ptx test.cu

結果を適当に抜粋。

        ld.global.f64   %fd1, [%r6+0];
        mov.f64         %fd2, 0d4014000000000000;       // 5
        mul.f64         %fd3, %fd1, %fd2;
        st.global.f64   [%r8+0], %fd3;

%r6には入力のアドレス、%r8には出力のアドレスが入っている。まあ妥当なコードか。次に実行コードCUBINを確認する。次のコマンドを打って、さらにdecudaを使って逆アセンブルする。

nvcc -arch sm_13 -O3 --cubin test.cu

結果を抜粋。

000028: 10008011 00000003 mov.b32 $r4, 0x00000000
000030: 10008015 04014003 mov.b32 $r5, 0x40140000
000038: d00e0209 80800780 mov.b64 $r2, g[$r1]
000040: e0040409 80000780 mul.rn.f64 $r2, $r2, $r4
000050: d00e0009 a0800781 mov.end.b64 g[$r0], $r2

これをみると次のことが推測できる。

  • ストリーム・プロセッサは32ビットのレジスタを持つので、double(64ビット)の値は、2つのレジスタに格納される。
  • double定数のレジスタへの代入は、32ビットmove命令を二回行うことで達成される。
  • doubleの演算命令のオペランドは3つ(出力レジスタ、入力レジスタ1、入力レジスタ2)だけなので、隣り合うレジスタ(例:$r4と$r5)に一つのdoubleを格納し、先頭のほうのレジスタを演算時に指定する(と推測される)。

ベンチマーク

デバイス情報

  • GeForce 8400 GS
    $ ./deviceQuery
    CUDA Device Query (Runtime API) version (CUDART static linking)
    There is 1 device supporting CUDA
    Device 0: "GeForce 8400 GS"
      CUDA Driver Version:                           2.30
      CUDA Runtime Version:                          2.30
      CUDA Capability Major revision number:         1
      CUDA Capability Minor revision number:         1
      Total amount of global memory:                 267714560 bytes
      Number of multiprocessors:                     1
      Number of cores:                               8
      Total amount of constant memory:               65536 bytes
      Total amount of shared memory per block:       16384 bytes
      Total number of registers available per block: 8192
      Warp size:                                     32
      Maximum number of threads per block:           512
      Maximum sizes of each dimension of a block:    512 x 512 x 64
      Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
      Maximum memory pitch:                          262144 bytes
      Texture alignment:                             256 bytes
      Clock rate:                                    1.40 GHz
      Concurrent copy and execution:                 No
      Run time limit on kernels:                     No
      Integrated:                                    No
      Support host page-locked memory mapping:       No
      Compute mode:                                  Default (multiple host threads can use this device simultaneously)
    
  • Tesla C1060
    $ ./deviceQuery
    CUDA Device Query (Runtime API) version (CUDART static linking)
    There are 2 devices supporting CUDA
    Device 0: "Tesla C1060"
      CUDA Driver Version:                           2.30
      CUDA Runtime Version:                          2.30
      CUDA Capability Major revision number:         1
      CUDA Capability Minor revision number:         3
      Total amount of global memory:                 4294705152 bytes
      Number of multiprocessors:                     30
      Number of cores:                               240
      Total amount of constant memory:               65536 bytes
      Total amount of shared memory per block:       16384 bytes
      Total number of registers available per block: 16384
      Warp size:                                     32
      Maximum number of threads per block:           512
      Maximum sizes of each dimension of a block:    512 x 512 x 64
      Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
      Maximum memory pitch:                          262144 bytes
      Texture alignment:                             256 bytes
      Clock rate:                                    1.30 GHz
      Concurrent copy and execution:                 Yes
      Run time limit on kernels:                     No
      Integrated:                                    No
      Support host page-locked memory mapping:       Yes
      Compute mode:                                  Default (multiple host threads can use this device simultaneously)
    
  • Number of multiprocessorsは、ストリーミングマルチプロセッサ(SM)の数
  • Number of coresは、ストリーミングプロセッサ(SP)の数(= SMの数 * 8)

バンド幅

  • GeForce 8400 GS
    $ ./bandwidthTest
    Running on......
          device 0:GeForce 8400 GS
    Quick Mode
    Host to Device Bandwidth for Pageable memory .
    Transfer Size (Bytes)   Bandwidth(MB/s)
     33554432               1926.5
    Quick Mode
    Device to Host Bandwidth for Pageable memory
    .
    Transfer Size (Bytes)   Bandwidth(MB/s)
     33554432               1396.0
    Quick Mode
    Device to Device Bandwidth
    .
    Transfer Size (Bytes)   Bandwidth(MB/s)
     33554432               4050.7
    
  • Tesla C1060
    $ ./bandwidthTest
    Running on......
          device 0:Tesla C1060
    Quick Mode
    Host to Device Bandwidth for Pageable memory
    .
    Transfer Size (Bytes)   Bandwidth(MB/s)
     33554432               4687.9
    Quick Mode
    Device to Host Bandwidth for Pageable memory
    .
    Transfer Size (Bytes)   Bandwidth(MB/s)
     33554432               3680.1
    Quick Mode
    Device to Device Bandwidth
    .
    Transfer Size (Bytes)   Bandwidth(MB/s)
     33554432               73361.7