見出し画像

CUDA プログラミング入門

旅人

この記事は、NVIDIA社が提供するCUDAプログラミングに関する教材の日本語訳です。

この記事は、NVIDIAの人気のある並列計算プラットフォームおよびプログラミングモデルであるCUDAの超簡単な入門編です。私は以前、「2013年に簡単なCUDA入門」という記事を書きましたが、それは長年にわたり人気がありました。しかし、CUDAプログラミングはますます簡単になり、GPUもはるかに高速化しているので、更新された（さらに簡単な）入門編を提供する時が来ました。

CUDA C++は、CUDAを使用して大規模な並列アプリケーションを作成する方法の一つです。強力なC++プログラミング言語を使用して、GPU上で実行される数千の並列スレッドによって加速された高性能アルゴリズムを開発することができます。多くの開発者がこの方法で計算と帯域幅を大量に消費するアプリケーションを加速しており、人工知能の革命を支えるライブラリやフレームワークにも使用されています。

CUDAについて聞いたことがあり、自分のアプリケーションでの使用に興味がある場合、CまたはC++プログラマーであれば、このブログ投稿が良いスタートを切る手助けになるでしょう。以下に従うためには、CUDA対応のGPUを搭載したコンピュータ（Windows、Mac、Linuxのいずれかで、NVIDIA GPUならどれでも）またはGPUを備えたクラウドインスタンス（AWS、Azure、IBM SoftLayerなどのクラウドサービスプロバイダーが提供しています）が必要です。無料のCUDA Toolkitもインストールしておく必要があります。また、クラウド上のGPUで実行されているJupyter Notebookを使用しても同様に従うことができます。

さあ、始めましょう！

始める

まず、100万要素の2つの配列の要素を加算する単純なC++プログラムから始めます。

#include <iostream>
#include <math.h>

// 2つの配列の要素を加算する関数
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 100万要素

  float *x = new float[N];
  float *y = new float[N];

  // ホストでxとy配列を初期化
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // CPUで100万要素を処理
  add(N, x, y);

  // エラーチェック（全ての値が3.0fであるべき）
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // メモリを解放
  delete [] x;
  delete [] y;

  return 0;
}

まず、このC++プログラムをコンパイルして実行します。上記のコードをファイルに保存して「add.cpp」として保存し、C++コンパイラでコンパイルします。私はMacを使用しているのでclang++を使用していますが、Linuxではg++、WindowsではMSVCを使用できます。

clang++ add.cpp -o add

次に実行します：

./add

期待通り、加算にエラーがなかったことを表示して終了します。次に、この計算をGPUの多数のコアで並列に実行したいと思います。最初のステップを踏むのは実際には非常に簡単です。

まず、`add`関数をGPUで実行できる関数に変える必要があります。これはCUDAでカーネルと呼ばれます。これを行うためには、関数に`global`指定子を追加するだけで、CUDA C++コンパイラにこの関数がGPUで実行され、CPUコードから呼び出すことができることを知らせます。

// GPUで2つの配列の要素を加算するCUDAカーネル関数
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

これらの`global`関数はカーネルとして知られ、GPUで実行されるコードはしばしばデバイスコードと呼ばれ、CPUで実行されるコードはホストコードと呼ばれます。

CUDAにおけるメモリ割り当て

GPUで計算するためには、GPUがアクセス可能なメモリを割り当てる必要があります。CUDAの統合メモリを使用すると、システム内のすべてのGPUおよびCPUがアクセスできる単一のメモリ空間を提供するため、これが簡単になります。統合メモリにデータを割り当てるには、`cudaMallocManaged()`を呼び出し、ホスト（CPU）コードまたはデバイス（GPU）コードからアクセスできるポインタを返します。データを解放するには、ポインタを`cudaFree()`に渡すだけです。

上記のコードで`new`の呼び出しを`cudaMallocManaged()`の呼び出しに置き換え、`delete []`の呼び出しを`cudaFree`の呼び出しに置き換える必要があります。

// 統合メモリを割り当てる - CPUまたはGPUからアクセス可能
float *x, *y;
cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));

...

// メモリを解放
cudaFree(x);
cudaFree(y);

最後に、`add()`カーネルを起動して、GPUで実行する必要があります。CUDAカーネルの起動は、三重角括弧構文`<<< >>>`を使用して指定されます。引数リストの前にこれを追加するだけです。

add<<<1, 1>>>(N, x, y);

簡単ですね！角括弧の中に何が入るかの詳細については後ほど説明しますが、今のところ、この行は`add()`を実行する1つのGPUスレッドを起動することを意味することだけを知っておけば十分です。

もう1つだけ必要なことがあります。カーネルが完了するまでCPUが結果にアクセスしないように待機させる必要があります（CUDAカーネルの起動は呼び出し元のCPUスレッドをブロックしないためです）。これを行うには、最終的なエラーチェックを行う前に`cudaDeviceSynchronize()`を呼び出すだけです。

完全なコードは以下の通りです：

#include <iostream>
#include <math.h>
// 2つの配列の要素を加算するカーネル関数
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y;

  // 統合メモリを割り当てる - CPUまたはGPUからアクセス可能
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // ホストでxとy配列を初期化
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // GPUで100万要素を処理
  add<<<1, 1>>>(N, x, y);

  // ホストでアクセスする前にGPUの終了を待つ
  cudaDeviceSynchronize();

  // エラーチェック（全ての値が3.0fであるべき）
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std

ここでは、NVIDIAの人気のある並列計算プラットフォームおよびプログラミングモデルであるCUDAの超簡単な紹介を行います。私は以前、2013年に「Easy Introduction to CUDA」という人気の投稿を書きました。しかし、CUDAプログラミングはより簡単になり、GPUもはるかに高速になったので、更新された（そしてさらに簡単な）紹介をする時が来ました。

CUDA C++は、CUDAを使用して大規模な並列アプリケーションを作成する方法の一つです。これは、強力なC++プログラミング言語を使用して、GPU上で動作する数千の並列スレッドによって加速される高性能アルゴリズムを開発することができます。多くの開発者が、深層学習として知られる人工知能の進行中の革命の基盤となるライブラリやフレームワークを含む、計算および帯域幅を大量に消費するアプリケーションをこの方法で加速してきました。

あなたはCUDAについて聞いたことがあり、あなた自身のアプリケーションで使用する方法を学びたいと思っているかもしれません。CまたはC++プログラマーであれば、このブログ投稿は良いスタートを提供するはずです。これに従うには、CUDA対応のGPUを備えたコンピュータ（Windows、Mac、またはLinux、そしてNVIDIAのGPUならどれでもよい）またはGPUを備えたクラウドインスタンス（AWS、Azure、IBM SoftLayer、その他のクラウドサービスプロバイダーが提供）と、無料のCUDAツールキットが必要です。また、クラウド上で動作するGPUを搭載したJupyter Notebookでも従うことができます。

さあ、始めましょう！

### シンプルなスタート
まず、100万要素の2つの配列の要素を加算する簡単なC++プログラムから始めます。

```cpp
#include <iostream>
#include <math.h>

// 2つの配列の要素を加算する関数
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 100万要素

  float *x = new float[N];
  float *y = new float[N];

  // ホスト上でxおよびy配列を初期化
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // CPU上で100万要素のカーネルを実行
  add(N, x, y);

  // エラーをチェック（すべての値が3.0fであるべき）
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // メモリを解放
  delete [] x;
  delete [] y;

  return 0;
}

まず、このC++プログラムをコンパイルして実行します。上記のコードをファイルに保存し、add.cppとして保存してから、C++コンパイラでコンパイルします。私はMacを使用しているのでclang++を使用しますが、Linuxではg++、WindowsではMSVCを使用できます。

> clang++ add.cpp -o add

次に実行します。

> ./add
 Max error: 0.000000

（Windowsでは実行可能ファイルをadd.exeと名付け、.\addで実行することをお勧めします。）

予想通り、加算にエラーがなかったことが表示され、終了します。次に、この計算をGPUの多くのコアで並列に実行したいと思います。実際には、最初のステップを踏むのは非常に簡単です。

まず、add関数をGPUが実行できる関数（CUDAではカーネルと呼ばれます）に変更するだけです。これを行うには、関数に__global__指定子を追加するだけで、これによりCUDA C++コンパイラがこれはGPU上で実行される関数であり、CPUコードから呼び出すことができることを認識します。

// GPU上で2つの配列の要素を加算するCUDAカーネル関数
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

これらの__global__関数はカーネルとして知られており、GPU上で実行されるコードはデバイスコードと呼ばれることが多く、CPU上で実行されるコードはホストコードと呼ばれます。

CUDAでのメモリ割り当て

GPUで計算を行うためには、GPUがアクセスできるメモリを割り当てる必要があります。CUDAの統一メモリを使用すると、システム内のすべてのGPUとCPUがアクセスできる単一のメモリ空間を提供するため、これが容易になります。統一メモリにデータを割り当てるには、cudaMallocManaged()を呼び出し、ホスト（CPU）コードまたはデバイス（GPU）コードからアクセスできるポインタを返します。データを解放するには、ポインタをcudaFree()に渡すだけです。

上記のコードのnew呼び出しをcudaMallocManaged()の呼び出しに置き換え、delete []の呼び出しをcudaFreeに置き換えます。

// 統一メモリの割り当て - CPUまたはGPUからアクセス可能
float *x, *y;
cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));

...

// メモリを解放
cudaFree(x);
cudaFree(y);

最後に、add()カーネルを起動する必要があります。これにより、GPU上で実行されます。CUDAカーネルの起動は三重角括弧構文<<< >>>を使用して指定されます。これをパラメータリストの前に追加するだけです。

add<<<1, 1>>>(N, x, y);

簡単ですね！ 角括弧の中に何が入るのかの詳細についてはすぐに説明しますが、今のところ、この行はadd()を実行するために1つのGPUスレッドを起動することだけを知っておけば十分です。

もう一つだけ。CPUがカーネルの結果にアクセスする前に、カーネルが完了するまで待つ必要があります（CUDAカーネルの起動は呼び出し元のCPUスレッドをブロックしないため）。これを行うには、cudaDeviceSynchronize()を呼び出して、CPUでの最終エラーチェックの前に呼び出します。

以下が完全なコードです：

#include <iostream>
#include <math.h>
// カーネル関数 - 2つの配列の要素を加算
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y;

  // 統一メモリの割り当て - CPUまたはGPUからアクセス可能
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // ホスト上でxおよびy配列を初期化
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // GPU上で100万要素のカーネルを実行
  add<<<1, 1>>>(N, x, y);

  // ホストでアクセスする前にGPUの完了を待つ
  cudaDeviceSynchronize();

  // エラーをチェック（すべての値が3.0fであるべき）
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(max

次に、CUDAカーネルを起動して並列に計算を行う方法について説明します。

### カーネル起動と並列処理
カーネルの`add<<<1, 1>>>(N, x, y)`という起動コードでは、`<<<1, 1>>>`という構文を使用しています。この部分を変更して、より多くのスレッドを使って並列に計算を行うようにします。

まず、GPUの計算は「グリッド」と「ブロック」の2つのレベルで構成されます。各グリッドは複数のブロックに分かれ、各ブロックは複数のスレッドを持ちます。これにより、膨大な数のスレッドが並列に計算を実行できます。

```cpp
__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

上記のカーネルでは、各スレッドが特定のインデックスを計算するように変更されました。これにより、複数のスレッドが協力して配列全体を処理できます。

カーネル起動の改善

カーネルを起動するとき、ブロックとスレッドの数を指定する必要があります。例えば、次のようにします。

int blockSize = 256; // 1ブロックあたりのスレッド数
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);

このコードは、256スレッドのブロックを使用し、必要なだけ多くのブロックを起動します。

完全なコード

以下が、すべての変更を含む最終的なコードです。

#include <iostream>
#include <math.h>

// GPU上で2つの配列の要素を加算するCUDAカーネル
__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 1百万要素
  float *x, *y;

  // 統合メモリの割り当て
  cudaMallocManaged(&x, N * sizeof(float));
  cudaMallocManaged(&y, N * sizeof(float));

  // 配列の初期化
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // カーネル起動のためのブロック数とスレッド数の設定
  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<<numBlocks, blockSize>>>(N, x, y);

  // カーネルの完了を待機
  cudaDeviceSynchronize();

  // エラーチェック
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i] - 3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // メモリの解放
  cudaFree(x);
  cudaFree(y);

  return 0;
}

このコードは、GPU上での並列処理の基本を理解するための素晴らしい出発点です。CUDAのより高度な機能について学びたい場合、公式のCUDAプログラミングガイドを参照することをお勧めします。

参考文献：

ーーーーーーーーーーーーーーーーーーーーーーーーーーーーーーーーーー

Twitterでぜひご意見をお寄せ下さい。フォローよろしくお願いします🙇

旅人Twitter
https://twitter.com/Tomoto1234567


この記事が気に入ったらサポートをしてみませんか？