【CUDA】CUDA入門、基本的なコーディング方法を学ぶ

CUDA入門
マウスコンピューター/G-Tune

CUDAのコーディング方法と基礎的な知識を解説していきます。
この記事を理解すれば、簡単な処理なら作れるようになるでしょう。

CUDAでは様々なAPIがありますが、すべてを把握してからではなく、まずは動く処理を作ることを目指しましょう。

基本的な知識

コーディングをするときに覚えておくべき知識を解説します。

用語

解説やコード中に出てくる用語は以下の通りです。
最低限、これらの用語を把握しておかないとCUDAの学習がままならないのでしっかり覚えておきましょう。

用語解説
ホストCPU側を指す。
ホストメモリCPU側のメモリ。
デバイスGPU側を指す。
デバイスメモリGPU側のメモリ。
基本的にホストメモリー>デバイスメモリにデータを転送してから処理を行う。
カーネルGPUの並列処理を記述する関数。

処理の流れ

CPUとGPUの処理の流れはこのようになります。

基本的にCPU側がGPU側に指示をしない限り、勝手にGPU側が動くことはありません

グリッドとブロックについて

カーネル実行時、どれだけのスレッド数で動かすかを指定する必要があります。

図の通り、スレッドはブロック内に複数あり、ブロックはグリッド内に複数あります。
また、ブロック数とスレッド数は3次元(x,y,z)で指定することができます。

ブロック数とスレッド数は以下のように「dim3」型で定義し、カーネル関数の<<< >>>で囲われた箇所に指定します。

と指定すると、駆動するスレッド数は x = 8 個, y = 1個, z = 1個 になります。 

スレッドをたくさん使用する場合、32の倍数が効率が良くなります。
さらにその中で、128個や256個が効率が上がりやすいようです。

コーディング

作る処理

作る処理の内容ですが、2つの配列の各要素を加算して、別の配列に代入していくという処理です。
各要素の加算はすべて別のスレッドで並列に行います。

API

今回のコーディングで使用するCUDAのAPIは以下の通り。

API名解説
cudaMallocHost/
cudaFreeHost
ホストメモリを確保/開放する。
cudaMalloc/
cudaFree
デバイスメモリを確保/開放する。
cudaMemcpyホストメモリとデバイスメモリ間でデータを転送する。

コーディング

コードは「.cu/.cuh」というファイルに記述していきます。

解説では、プロジェクトに「CudaProcess.cu」 「CudaProcess.cuh」というファイルを追加して記述しています。

Process() という関数に一連の処理を記述しています。

メモリの確保->データ転送->カーネル実行->結果転送->メモリ開放
という流れです。

カーネル内で
[int idx = blockDim.x * blockIdx.x + threadIdx.x;]
と記述している箇所では、自分を呼び出しているスレッドのindexを取得しています。
よく使う記述なので覚えておきましょう。

blockDim :ブロック内のスレッド数
blockIdx : ブロックIndex
threadIdx : スレッドのIndex

注意点

  • ホストメモリはGPU側からはアクセスできません。
    (例外的に、ホストメモリに直接 アクセスしているっぽく作る方法はある)
  • デバイスメモリもCPU側からはアクセスできません。
  • カーネルは戻り値を持てません。結果はデバイスメモリに書き込んで、ホストメモリへ転送する必要があります。
  • 解説中では、メモリ確保からカーネル実行までをまとめてやっています。しかし、メモリの確保は時間がかかるので、本来はソフト起動時などにあらかじめやっておくほうが良いです。
  • スレッド数の最大は x=1024個 y=1024個 z=64個 までです。

補足

以下の記事では、追加でGPUの特徴や用語を解説しています。

【CUDA】Grid,Block,Thread,Warpについて
Grid,Block,Thread,Warp(グリッド、ブロック、スレッド、ワープ)の関係について解説していきます。

以下の記事では、ブロックやスレッドを2次元で扱ったときのコード例を紹介しています。

【CUDA】ブロック・スレッドを2次元で使う
CUDAのスレッドは2次元で持つことができるので、その使い方の解説をします。
マウスコンピューター/G-Tune