cudacudaを使い始める


備考

CUDAは、GPUのための独自のNVIDIA並列コンピューティング技術とプログラミング言語です。

GPUは、並行して何千もの軽量スレッドを実行できる高度な並列マシンです。通常、各GPUスレッドは実行速度が遅く、コンテキストが小さくなります。一方、GPUは数千のスレッドを並列に、さらには同時に実行することができます(正確な数値は実際のGPUモデルに依存します)。 CUDAは、NVIDIA GPUアーキテクチャ専用に設計されたC ++の方言です。しかし、アーキテクチャの相違により、ほとんどのアルゴリズムは単純なC ++から単純にコピー&ペーストすることはできません。実行されますが、非常に遅くなります。

用語

  • ホスト - 通常のCPUベースのハードウェアとその環境で動作する通常のプログラムを指します
  • デバイス - CUDAプログラムが実行する特定のGPUを指します。単一のホストが複数のデバイスをサポートできます。
  • カーネル - ホストコードから呼び出せるデバイスに常駐する関数。

物理プロセッサ構造

CUDA対応GPUプロセッサの物理構造は次のとおりです。

  • チップ - GPUのプロセッサ全体。いくつかのGPUには2つのGPUがあります。
  • ストリーミングマルチプロセッサ (SM) - 各チップには、モデルに応じて〜100個までのSMが含まれています。各SMは、互いに独立して動作し、グローバルメモリのみを使用して相互に通信します。
  • CUDAコア - SMの単一のスカラー計算ユニット。正確な数はアーキテクチャによって異なります。各コアは、(CPUのハイパースレッディングと同様に)素早く連続して同時に実行されるいくつかのスレッドを処理できます。

さらに、各SMは、1つ以上のワープスケジューラを特徴とする。各スケジューラは、1つの命令をいくつかのCUDAコアにディスパッチします。これにより、効果的にSMが32ワイドSIMDモードで動作します。

CUDA実行モデル

GPUの物理的構造は、カーネルがデバイス上でどのように実行され、どのようにCUDAでそれらをプログラミングするかに直接影響します。カーネルは、呼び出される並列スレッドの数を指定する呼び出し構成で呼び出されます

  • グリッドは、カーネル呼び出し時に生成されるすべてのスレッドを表します。これは、 ブロックの 1つまたは2つの次元的なセットとして指定されます
  • ブロック - は半独立したスレッドセットです。各ブロックは単一のSMに割り当てられます。そのため、ブロックはグローバルメモリを介してのみ通信できます。ブロックは決して同期されません。あまりにも多くのブロックがある場合、一部は他のブロックの後で連続して実行することができます。一方、リソースが許可されている場合、同じSM上で複数のブロックが実行される可能性がありますが、プログラマはそれが恩恵を受けることはできません(明らかなパフォーマンスの向上を除く)。
  • スレッド - 単一のCUDAコアによって実行されるスカラーシーケンスの命令。スレッドはコンテキストを最小限にした「軽量」なので、ハードウェアを素早く入れ替えることができます。その数のために、CUDAスレッドは、それらに割り当てられた少数のレジスタと非常に短いスタックで動作します(できれば全くありません!)。そのため、CUDAコンパイラは、静的なジャンプとループのみを含むようにカーネルをフラット化するために、すべての関数呼び出しをインライン化することを推奨します。多くの新しいデバイスでサポートされている間に、関数呼び出し呼び出しと仮想メソッド呼び出しは、通常、大きなパフォーマンス上のペナルティを被ります。

各スレッドは、 threadIdxブロック内のブロックインデックスblockIdxとスレッドインデックスによって識別されます。これらの数値は、実行中のスレッドによっていつでもチェックすることができ、スレッドを別のスレッドと区別する唯一の方法です。

さらに、スレッドは、それぞれが正確に32のスレッドを含むワープに編成されています。単一のワープ内のスレッドは、SIMDファシオンで完璧な同期で実行されます。異なるワープからのスレッドは同じブロック内で任意の順序で実行できますが、プログラマによって強制的に同期させることができます。異なるブロックからのスレッドは、どのような方法でも直接同期または相互作用することはできません。

メモリ構成

通常のCPUプログラミングでは、メモリ構成は通常プログラマから隠されています。典型的なプログラムは、ちょうどRAMがあるかのように動作します。レジスタ管理、L1- L2- L3-キャッシング、ディスクへのスワッピングなどのすべてのメモリ操作は、コンパイラ、オペレーティングシステム、またはハードウェア自体によって処理されます。

これはCUDAのケースではありません。より新しいGPUモデルは、例えばCUDA 6のUnified Memoryなどの部分的な負担を隠していますが、パフォーマンス上の理由から組織を理解する価値はあります。基本的なCUDAメモリ構造は次のとおりです。

  • ホストメモリ - 通常のRAM主にホストコードで使用されますが、新しいGPUモデルでも同様にアクセスできます。カーネルがホストメモリにアクセスするとき、GPUは通常、PCIeコネクタを介してマザーボードと通信する必要があり、そのため比較的遅いです。
  • デバイスメモリ/グローバルメモリ - GPUのメインメモリで、すべてのスレッドが使用できます。
  • 共有メモリ - 各SMに配置されているため、グローバルよりもはるかに高速にアクセスできます。共有メモリは各ブロックに専用です。 1つのブロック内のスレッドは、それを通信に使用できます。
  • レジスタ - 各スレッドの最も速く、プライベートな、アドレス不定のメモリ。一般に、これらは通信に使用することはできませんが、いくつかの組み込み関数では、その内容をワープ内でシャッフルすることができます。
  • ローカルメモリ -アドレス指定可能である各スレッドのプライベートメモリ。これは、レジスタ流出、および可変インデックスを持つローカル配列に使用されます。物理的には、それらはグローバルメモリに存在します。
  • テクスチャメモリ、定数メモリ - グローバルメモリの一部で、カーネルに対して不変であるとマークされています。これにより、GPUは専用キャッシュを使用できます。
  • L2キャッシュ - オンチップで、すべてのスレッドが使用できます。スレッドの量を考えると、各キャッシュラインの予想寿命はCPUよりもはるかに短くなります。これは主に、ミスアラインと部分的にランダムなメモリアクセスパターンを支援するために使用されます。
  • L1キャッシュ - 共有メモリと同じスペースにあります。ここでも、量はそれを使用するスレッドの数を考えるとやや小さいので、データがそこに長時間留まることは期待しないでください。 L1キャッシングを無効にすることができます。

バージョン

コンピューティング能力建築 GPUコード名発売日
1.0 テスラ G80 2006年11月8日
1.1 テスラ G84、G86、G92、G94、G96、G98、 2007-04-17
1.2 テスラ GT218、GT216、GT215 2009-04-01
1.3 テスラ GT200、GT200b 2009-04-09
2.0 フェルミ GF100、GF110 2010-03-26
2.1 フェルミ GF104、GF106 GF108、GF114、GF116、GF117、GF119 2010-07-12
3.0 ケプラー GK104、GK106、GK107 2012-03-22
3.2 ケプラー GK20A 2014-04-01
3.5 ケプラー GK110、GK208 2013-02-19
3.7 ケプラー GK210 2014-11-17
5.0 マクスウェル GM107、GM108 2014-02-18
5.2 マクスウェル GM200、GM204、GM206 2014-09-18
5.3 マクスウェル GM20B 2015-04-01
6.0 パスカル GP100 2016-10-01
6.1 パスカル GP102、GP104、GP106 2016年5月27日

リリース日は、与えられたコンピューティング能力をサポートする最初のGPUのリリースとなります。いくつかの日付はおおよそのものです。たとえば、2014年第2四半期に3.2カードがリリースされました。

サンプルプログラムのコンパイルと実行

NVIDIAインストールガイドはサンプルプログラムを実行して終了し、CUDA Toolkitのインストールを確認しますが、明示的に記載していません。まず、すべての前提条件を確認します。サンプル・プログラムのデフォルトのCUDAディレクトリーを確認してください。存在しない場合は、公式のCUDAウェブサイトからダウンロードできます。例が存在するディレクトリに移動します。

$ cd /path/to/samples/
$ ls
 

次のような出力が表示されます。

0_Simple     2_Graphics  4_Finance      6_Advanced       bin     EULA.txt
1_Utilities  3_Imaging   5_Simulations  7_CUDALibraries  common  Makefile
 

このディレクトリにMakefile が存在することを確認してください。 UNIXベースのシステムでmake コマンドを実行すると、すべてのサンプルプログラムがビルドされます。または、別のMakefile が存在するサブディレクトリに移動し、そこからmake コマンドを実行しmake そのサンプルのみをビルドします。

2つの推奨サンプルプログラム、 deviceQuerybandwidthTest 実行します。

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 
 

出力は次のようになります。

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 950M"
  CUDA Driver Version / Runtime Version          7.5 / 7.5
  CUDA Capability Major/Minor version number:    5.0
  Total amount of global memory:                 4096 MBytes (4294836224 bytes)
  ( 5) Multiprocessors, (128) CUDA Cores/MP:     640 CUDA Cores
  GPU Max Clock rate:                            1124 MHz (1.12 GHz)
  Memory Clock rate:                             900 Mhz
  Memory Bus Width:                              128-bit
  L2 Cache Size:                                 2097152 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 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 7.5, CUDA Runtime Version = 7.5, NumDevs = 1, Device0 = GeForce GTX 950M
Result = PASS
 

最後のResult = PASS というステートメントは、すべてが正しく機能していることを示します。今度は、他のサンプルプログラムのbandwidthTest を同様の方法で実行します。出力は次のようになります。

[CUDA Bandwidth Test] - Starting...
Running on...

 Device 0: GeForce GTX 950M
 Quick Mode

 Host to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10604.5

 Device to Host Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10202.0

 Device to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            23389.7

Result = PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
 

ここでも、 Result = PASS 文はすべてが正しく実行されたことを示します。他のすべてのサンプルプログラムも同様の方法で実行できます。

CUDAスレッドを1つ起動して、こんにちは

このシンプルなCUDAプログラムは、GPU(別名「デバイス」)上で実行される関数を記述する方法を示しています。 CPU(「ホスト」)は、「カーネル」と呼ばれる特別な関数を呼び出してCUDAスレッドを作成します。 CUDAプログラムは、構文が追加されたC ++プログラムです。

どのように動作するかを確認するには、 hello.cu という名前のファイルに次のコードを記述します。

#include <stdio.h>

// __global__ functions, or "kernels", execute on the device
__global__ void hello_kernel(void)
{
  printf("Hello, world from the device!\n");
}

int main(void)
{
  // greet from the host
  printf("Hello, world from the host!\n");

  // launch a kernel with a single thread to greet from the device
  hello_kernel<<<1,1>>>();

  // wait for the device to finish so that we see the message
  cudaDeviceSynchronize();

  return 0;
}
 

(デバイス上でprintf 関数を使用するには、少なくとも2.0の計算能力を持つデバイスが必要です(詳細については、 バージョンの概要を参照してください)。

今度はNVIDIAコンパイラを使ってプログラムをコンパイルして実行しましょう:

$ nvcc hello.cu -o hello
$ ./hello
Hello, world from the host!
Hello, world from the device!
 

上記の例に関する追加情報

  • nvcc は "NVIDIA CUDA Compiler"の略です。これは、ソースコードをホストコンポーネントとデバイスコンポーネントに分離します。
  • __global__ は、関数がGPUデバイス上で実行され、ホストから呼び出されたことを示す関数宣言で使用されるCUDAキーワードです。
  • <<< 括弧( <<<>>> )は、ホストコードからデバイスコード(「カーネル起動」とも呼ばれます)への呼び出しをマークします。これらの三角括弧内の数字は、並列実行する回数とスレッド数を示します。

前提条件

CUDAでプログラミングを開始するには、 CUDA Toolkitと開発者用ドライバをダウンロードしてインストールします。このツールキットには、 nvcc 、NVIDIA CUDA Compiler、およびCUDAアプリケーション開発に必要なその他のソフトウェアが含まれています。ドライバは、GPUプログラムがCUDA対応ハードウェア上で正しく動作することを保証します。

コマンドラインからnvcc --version を実行すると、CUDA Toolkitがマシンに正しくインストールされているかどうかを確認できます。たとえば、Linuxマシンでは、

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Tue_Jul_12_18:28:38_CDT_2016
Cuda compilation tools, release 8.0, V8.0.32
 

コンパイラ情報を出力します。前のコマンドが成功しなかった場合は、CUDAツールキットがインストールされていないか、またはnvcc (Windowsマシンの場合はC:\CUDA\bin /usr/local/cuda/bin 、POSIX OSの場合は/usr/local/cuda/bin への/usr/local/cuda/binPATH 環境変数。

さらに、CUDAプログラムをコンパイルおよび構築するためにnvcc で動作するホストコンパイラも必要です。 Windowsでは、これはMicrosoft Visual Studioに同梱されているMicrosoftコンパイラのcl.exe です。 POSIX OSでは、 gccg++ を含む他のコンパイラが利用できます。公式のCUDA クイックスタートガイドでは、特定のプラットフォームでサポートされているコンパイラのバージョンを知ることができます。

すべてが正しく設定されていることを確認するには、すべてのツールが正しく動作するように、簡単なCUDAプログラムをコンパイルして実行してみましょう。

__global__ void foo() {}

int main()
{
  foo<<<1,1>>>();

  cudaDeviceSynchronize();
  printf("CUDA error: %s\n", cudaGetErrorString(cudaGetLastError()));

  return 0;
}
 

このプログラムをコンパイルするには、test.cuというファイルにコピーし、コマンドラインからコンパイルします。たとえば、Linuxシステムでは、次のように動作します。

$ nvcc test.cu -o test
$ ./test
CUDA error: no error
 

プログラムがエラーなく成功したら、コーディングを始めましょう!

CUDAで2つの配列を合計する

この例は、2つのint 配列をCUDAで合計する単純なプログラムを作成する方法を示しています。

CUDAプログラムは異種であり、CPUとGPUの両方で実行される部品で構成されています。

CUDAを利用するプログラムの主要部分は、CPUプログラムに似ており、

  • GPUで使用されるデータのメモリ割り当て
  • ホストメモリからGPUメモリへのデータコピー
  • カーネル関数を呼び出してデータを処理する
  • 結果をCPUメモリにコピーする

デバイスのメモリを割り当てるには、 cudaMalloc 関数を使用します。デバイスとホストの間でデータをコピーするには、 cudaMemcpy 関数を使用できます。 cudaMemcpy の最後の引数は、コピー操作の方向を指定します。可能なタイプは5つあります。

  • cudaMemcpyHostToHost - ホスト - >ホスト
  • cudaMemcpyHostToDevice - ホスト - >デバイス
  • cudaMemcpyDeviceToHost - デバイス - >ホスト
  • cudaMemcpyDeviceToDevice - デバイス - >デバイス
  • cudaMemcpyDefault - デフォルトベースの統一仮想アドレス空間

次に、カーネル関数が呼び出されます。トリプルシェブロン間の情報は実行コンフィギュレーションであり、並列にカーネルを実行するデバイススレッドの数を指定します。最初の数値(例では2 )はブロック数を指定し、 2 番目の数値((例では(size + 1) / 2 ) - ブロック内のスレッド数を指定します。この例では、1つのスレッドが2つの要素を担当するのではなく、1つの余分なスレッドを要求するように、サイズに1を追加することに注意してください。

カーネルの呼び出しは非同期関数なcudaDeviceSynchronize 、実行が完了するまで待機するためにcudaDeviceSynchronize が呼び出されます。結果配列はホストメモリにコピーされ、デバイスに割り当てられたすべてのメモリはcudaFree で解放されcudaFree

関数をカーネルとして定義するには、 __global__ 宣言指定子が使用されます。この関数は各スレッドによって呼び出されます。各スレッドが結果の配列の要素を処理するようにするには、各スレッドを識別して識別する手段が必要です。 CUDAは、変数blockDimblockIdx 、およびthreadIdx 定義します。事前定義された変数blockDim には、カーネルの起動のための2番目の実行コンフィギュレーションパラメータで指定された各スレッドブロックのディメンションが含まれます。事前定義された変数threadIdx およびblockIdx は、それぞれそのスレッドブロック内のスレッドおよびグリッド内のスレッドブロックのインデックスを含む。配列の要素よりもスレッドを1つ要求する可能性があるため、配列の最後を過ぎてアクセスしないように、 size を渡す必要があります。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void addKernel(int* c, const int* a, const int* b, int size) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < size) {
        c[i] = a[i] + b[i];
    }
}

// Helper function for using CUDA to add vectors in parallel.
void addWithCuda(int* c, const int* a, const int* b, int size) {
    int* dev_a = nullptr;
    int* dev_b = nullptr;
    int* dev_c = nullptr;

    // Allocate GPU buffers for three vectors (two input, one output)
    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaMalloc((void**)&dev_b, size * sizeof(int));

    // Copy input vectors from host memory to GPU buffers.
    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.
    // 2 is number of computational blocks and (size + 1) / 2 is a number of threads in a block
    addKernel<<<2, (size + 1) / 2>>>(dev_c, dev_a, dev_b, size);
    
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaDeviceSynchronize();

    // Copy output vector from GPU buffer to host memory.
    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
}

int main(int argc, char** argv) {
    const int arraySize = 5;
    const int a[arraySize] = {  1,  2,  3,  4,  5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    addWithCuda(c, a, b, arraySize);

    printf("{1, 2, 3, 4, 5} + {10, 20, 30, 40, 50} = {%d, %d, %d, %d, %d}\n", c[0], c[1], c[2], c[3], c[4]);

    cudaDeviceReset();

    return 0;
}