cuda チュートリアル
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/binはPATH環境変数。

さらに、CUDAプログラムをコンパイルおよび構築するためにnvccで動作するホストコンパイラも必要です。 Windowsでは、これはMicrosoft Visual Studioに同梱されているMicrosoftコンパイラのcl.exeです。 POSIX OSでは、 gccやg++を含む他のコンパイラが利用できます。公式の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は、変数blockDim 、 blockIdx 、および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;
}

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キーワードです。
• <<<括弧（ <<< 、 >>> ）は、ホストコードからデバイスコード（「カーネル起動」とも呼ばれます）への呼び出しをマークします。これらの三角括弧内の数字は、並列実行する回数とスレッド数を示します。

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

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つの推奨サンプルプログラム、 deviceQueryとbandwidthTest実行します。

$ 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文はすべてが正しく実行されたことを示します。他のすべてのサンプルプログラムも同様の方法で実行できます。