CUDAのインストールとサンプル実行

参考文献：Interface 2022年8月号　第２部第２章

環境構築

• WSL2上にAlmaLinuxをインストールする
• CUDA を Linux 上にインストールする

WSL の概要とインストール

Windows Subsystem for Linuxの略で、Windows の中でLinuxを動かすことができるようになります。

wsl　--install

>>> wsl --version
WSL バージョン: 2.6.1.0
カーネル バージョン: 6.6.87.2-1
WSLg バージョン: 1.0.66
MSRDC バージョン: 1.2.6353
Direct3D バージョン: 1.611.1-81528511
DXCore バージョン: 10.0.26100.1-240331-1435.ge-release
Windows バージョン: 10.0.26200.7171

AlmaLinux の概要とインストール

Linux ディストリビューションの一種のようです。
テキストに出てきた Ubuntu もその一種らしいです。

OSの中核部分を担うものをカーネルといい、カーネルであるLinuxとその他諸々のソフトウェアが一緒になったものをディストリビューション（; 分配 ; 配布）といいます。
これはまだインストールされていない（たぶん）ので、AlmaLinuxをインストールしていきます。

「Linux」はOSの名前として知られていますが、厳密にはカーネルの名前で、カーネルのLinux（狭義のLinux）は、その他諸々のソフトウェアと組み合わせることによって、OSのLinux（広義のLinux）として動きます。
https://wa3.i-3-i.info/word12473.html

wsl --install AlmaLinux-9

ユーザー名とパスワードの設定が求められました。パスワードはターミナルに入力しても表示されないのですが、何も表示されなくてもきちんと入力して Enter キーを押すと、通りました。謎です。

wsl と AlmaLinux のログインログアウト

• ディストリビューションと合わせて起動

wsl -d AlmaLinux-9

• 現在の AlmaLinux セッションを終了する

exit

• wsl 完全停止

wsl --shutdown

• wsl の状態を確認する

wsl -l -v

CUDA（クーダ） の概要とインストール

GPU（Graphics Processing Unit）をグラフィックの描画以外の計算などにも使えるようにするためのプラットフォームだそうです。

CUDA が提供されることによって、NVIDIA の GPU をプログラミングするためのモデルと言語が提供された（テキストから抜粋）

並列計算に便利な GPU を他の用途でも使おうよ！的な感じだと解釈しました。

以下のサイトに従ってインストールしました。
CUDA on WSL ユーザーガイド

途中インストールの仕方で Ubuntu 用のコードが出てきたのですが、それだとエラーがでたので、dnf-plugins-core をインストールしてから Toolkit を入れました。

# dnf config-manager インストール
sudo dnf install -y dnf-plugins-core

# NVIDIA CUDA リポジトリ追加
sudo dnf config-manager --add-repo https://developer.download.nvidia.com/compute/cuda/repos/rhel9/x86_64/cuda-rhel9.repo

sudo dnf clean all

# CUDA Toolkit 13.0 をインストール
sudo dnf -y install cuda-toolkit-13-0

NVIDIA グラフィクスドライバーと CUDA Toolkit が導入できました。

CUDA用 C/C++ コンパイラ NVCC

普通の C コンパイラとして.c のファイルを実行できる。（使用しているのはCPUのみ）
CUDAでは .cu という拡張子を使うと、CPU も GPU も使うことができるそうです。

Linux 上で環境変数を設定する

nvcc を見つけてもらうために、環境変数を設定します。

echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc

• echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc

  • /usr/local/cuda/bin を PATH に追加
  • これによりターミナルから nvcc などの CUDA コマンドを直接実行可能
  • >> ~/.bashrc で設定をホームディレクトリの .bashrc に追記

• echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc

  • CUDA のライブラリディレクトリ /usr/local/cuda/lib64 を LD_LIBRARY_PATH に追加
  • プログラムが CUDA のライブラリを見つけられるようにする
  • 同じく .bashrc に追記

• source ~/.bashrc

  • .bashrc に書いた設定を 今のターミナルに即時反映
  • 次回ターミナルを開いたときも設定が有効になる

Linux 上でファイルを作って開く

全然ここじゃない別の場所で編集できそうだが、一旦これで実行ファイルを作成しました。
pwd（コマンドプロンプトでいう cd）で /ユーザー名 にいることを確認して、以下のようにファイルを作成します。

nano hello.cu

nano の画面が出てくるので、そこに書き込みます。（C++ の Hello World）
C-o で保存、C-x で終了です。

#include <iostream>
int main()
{
    std::cout << "Hello CU World"
    << std::endl;
}

C と C++ の表示方法の違い

項目	C	C++
使うヘッダ	#include <stdio.h>	#include <iostream>
表示する関数 / 機能	printf()	std::cout
改行	\n	std::endl
出力の書き方	printf("Hello\n");	std::cout << "Hello" << std::endl;
言語の位置づけ	手続き型	オブジェクト指向対応

CUDA の基本言語は C++ です。

公式のサンプルで CUDA プログラムの構造を確認する

0.　Linux 上に git がない場合はインストールする
-y を入れると途中の y/N を y で確定して進めてくれます

sudo dnf install -y git

1.　公式サンプルのリポジトリが Github で公開されているのでクローンする

$ git clone https://github.com/NVIDIA/cuda-samples.git

2.　README に「CMakeを使用してビルドされています。」とあるので、CMake をインストールする

sudo dnf install -y cmake

3.　サンプルのあるディレクトリに移動して、ビルド用のフォルダを作る
ソースコードと、オブジェクトファイル（中間）・実行ファイル（機械語）がごちゃごちゃにならないように、サンプルごとにビルド用のディレクトリを作成します。
そして作ったビルド用ディレクトリに移動します。

cd cuda-samples/Samples/5_Domain_Specific/simpleVulkan
mkdir build
cd build

ビルドとサンプル分けのまとめ

用語	説明	例え
ソースコード	人間が書いたプログラム	材料
コンパイル	ソースコード → オブジェクトファイル（.o）	材料を下ごしらえ
リンク	複数のオブジェクトファイル → 実行ファイル	下ごしらえした材料を組み合わせて料理完成
ビルド	コンパイル + リンクの全作業	料理を作る一連の工程

4.　CMake を使ってビルド設定を行う
サンプルのソースコードがあるディレクトリを指定します。今回は一つ上

cmake ..

CUDA アーキテクチャが見つかりませんでしたというエラーは手動設定することで改善しました。wsl などの特殊環境下では起こりやすいみたいです。
一度 build を削除してもう一度作り直してから以下を実行しました。CUDA では GPU の世代ごとに最適なバイナリを作る必要があるそうで、世代によって指定する値は変わります。

cmake .. -DCMAKE_CUDA_ARCHITECTURES=75

5.　ビルドする

make

.o という拡張子のオブジェクトファイルや実行ファイルが作られます。

6.　実行する

./<サンプルの実行ファイル名>

サンプル : 5_Domain_Specific/simpleVulkan を実行する

simpleVulkanを実行してみます。
README より実行には x11 と Vulkan が必要とのことです。
x11 は Linux ディストリビューションに入ってるかも的なことが書いてありました。
また1度実行した時に、OpenGL と GLFW が見つかりませんというエラーも出たので、これらもインストールしていきます。

Vulkan とは

高性能・低オーバーヘッドな設計になっている描画処理を行うためのAPIだそうです。

sudo dnf install -y vulkan-loader-devel vulkan-tools

OpenGL とは

「ハードウェアに依存しない2D/3Dグラフィックス描画のための業界標準API」だそうで、
OpenGL を使うと、GPUに対する 絵や3Dグラフィックを書くための命令 が簡単に行えるそうです。

インストール例

sudo dnf install -y mesa-libGL mesa-libGL-devel

OpenGL プログラムの基本構造

ステップ	目的
① ウィンドウ作成	描画する画面を準備
② 描画関数作成	GPUに何を描くか指示
③ 描画ループ	描画を繰り返し更新
④ 終了処理	メモリやリソースを解放

GLFW とは

OpenGL APIを介してグラフィックス出力のみを行うアプリケーション用の、ウィンドウの作成と管理を可能にするライブラリのこと。
AlmaLinux は 最小構成なので GUI・グラフィックス系ライブラリは EPEL という追加リポジトリに分離されているようです。なので、EPAL を有効化してからインストールします。

# EPAL を有効化
sudo dnf install -y epel-release

# GLFW のインストール
sudo dnf install -y glfw glfw-devel

glfw3 not found というエラーがでました。調べてみると、CMake にライブラリの使い方を教えるための pkg-config がないため、glfw3 がないことにされたみたいです。

# pkg-config のインストール
sudo dnf install -y pkgconf-pkg-config

インストール後、以下を実行します。

cd ~/cuda-samples/Samples/5_Domain_Specific/simpleVulkan
mkdir build
cd build
cmake .. -DCMAKE_CUDA_ARCHITECTURES=75
make
./simpleVulkan

>>>
[ 20%] Building CXX object CMakeFiles/simpleVulkan.dir/main.cpp.o
[ 40%] Building CUDA object CMakeFiles/simpleVulkan.dir/SineWaveSimulation.cu.o
[ 60%] Building CXX object CMakeFiles/simpleVulkan.dir/VulkanBaseApp.cpp.o
[ 80%] Linking CUDA device code CMakeFiles/simpleVulkan.dir/cmake_device_link.o
[100%] Linking CXX executable simpleVulkan
[100%] Built target simpleVulkan

ビルドは無事成功しました。
調べたところ、wsl2 上では計算はできるが、表示まではできないとのことです。（No suitable device found! というエラーがでる）
（wslg が認識できているけど使われていない...? それとも使ったうえでできない...?）

サンプル : 5_Domain_Specific/nbody を実行する

GLUT が必要なので、インストールします。

sudo dnf install freeglut freeglut-devel mesa-libGL mesa-libGL-devel

実行します。

cd cuda-samples/Samples/5_Domain_Specific/nbody
mkdir build
cd build
cmake .. -DCMAKE_CUDA_ARCHITECTURES=75
make
./nbody

ビルドはできましたが、wsl2 上では表示できません。

> Compute 7.5 CUDA device: [NVIDIA T400 4GB]
CUDA error at /mnt/c/Users/nomura/cuda-samples/Samples/5_Domain_Specific/nbody/bodysystemcuda_impl.h:192 code=304(cudaErrorOperatingSystem) "cudaGraphicsGLRegisterBuffer(&m_pGRes[i], m_pbo[i], cudaGraphicsMapFlagsNone)"

サンプル : 0_Introduction/vectorAdd でプログラムの構成を理解する

cd cuda-samples/Samples/0_Introduction/vectorAdd
mkdir build
cd build
cmake .. -DCMAKE_CUDA_ARCHITECTURES=75
make
./vectorAdd

サンプル５の２つと違って、計算のみのコードなので、問題なく動きました。
テスト計算もパスできているようです。

>>>
Copy input data from the host memory to the CUDA device
CUDA kernel launch with 196 blocks of 256 threads
Copy output data from the CUDA device to the host memory
Test PASSED
Done

vectorAdd.cu ファイルを参照してみると、__global__ void vectorAdd() という関数と、int main(void){} の２つで構成されていました。

カーネル関数とは

• CUDA では GPU で作動するプログラムをカーネルと呼ぶ
• __global__ はカーネル関数であることを明示する
• この関数は CPU 側からも呼び出し可能である
• GPU で並列計算を行うコードなので、書かれてある内容は本来 for 文の中に書かれるような内容である

__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < numElements) {
        C[i] = A[i] + B[i] + 0.0f;
    }
}

for 文のように、何回繰り返すのかを示す条件が必要なので、引数に numElements があります。

カーネル関数を呼び出す側

int threadsPerBlock = 256;
int blocksPerGrid   = (numElements + threadsPerBlock - 1) / threadsPerBlock;
printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock);
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);

vectorAdd 関数についている、 <<<blocksPerGrid, threadsPerBlock>>> について、
CUDA におけるグリッド、ブロック、スレッズの関係性は以下の通りです。

GPU 全体 (Grid)
├─ Block 0 (グループ0)
│   ├─ Thread 0 (作業0)
│   ├─ Thread 1 (作業1)
│   └─ ...
├─ Block 1 (グループ1)
└─ Block 2

• blocksPerGrid：グリッド全体のブロック数
• threadsPerBlock：1ブロックあたりのスレッド数

例えば1000回計算を行う場合、T400のコア数384に合わせて割り振ると、処理の流れは
　Time 1: GPUコア0~383 にスレッド0~383が割り当てられる
　Time 2: GPUコア0~383 にスレッド384~767が割り当てられる
　Time 3: GPUコア0~383 にスレッド768~999が割り当てられる
のようになります。（実際は違うかもしれないけれど雰囲気）

並列処理を行うということはすなわち、スレッド0の作業を行ったコアが次に行う作業はスレッド1ではなく、コア数依存のスレッド i になるため、 i を毎回計算しなければならない。

そこでカーネル関数内で以下の計算を行う

int i = blockDim.x * blockIdx.x + threadIdx.x;

• blockDim.x : threadsPerBlock（ブロック内スレッド数）
• blockIdx.x : このブロックがグリッドの何番目か を表す
• threadIdx.x : ブロック内のスレッド番号

この計算によってすべてのスレッドが網羅できます。

今回この vectorAdd サンプルでは、１ブロックあたりのスレッド数は 256 でした。（blocksPerGrid は threadsPerBlock と 全スレッド数より計算で求めている）
NVIDIA GPU は 32 スレッド = 1 ワープ で同時に命令を実行する仕様で、ワープごとに順番に実行され、GPU は複数ワープをほぼ同時に処理しているそうです。
32 の倍数かつ、多すぎず少なすぎない 256 がよく使われるようです。

SM とは

SM は Streaming Multiprocessor の略で、CUDAコア、メモリ階層、ワープスケジューラ、演算ユニットなどを内包した実行処理を行う GPU の心臓部です。参考元

GPU全体
├── SM 1
│   ├── CUDA Core × N （GPUによる）
│   ├── 共有メモリ
│   ├── レジスタファイル
│   └── ワープスケジューラ × 4
├── SM 2
├── ...
└── SM M（GPUによる）

例えば 1024 個の計算を GPU で行う場合、

1. CPU → GPU 「256スレッドのブロックを4つ実行してください」
    

2. ブロックをSMに割り当てる
   このとき、ブロックは SM を跨がない。SMが６つだとすると

   SM0 ← Block0
   SM1 ← Block1
   SM2 ← Block2
   SM3 ← Block3
   SM4 ← 空き
   SM5 ← 空き
    

3. 各 SM が与えられたブロックをワープ単位で分解する

   Block0
   ├ Warp0 (thread 0–31)
   ├ Warp1 (32–63)
   ├ ...
   └ Warp7 (224–255)
    

4. 各 SM が 自身に割り当てられたワープの中から実行可能なワープを実行する
   （メモリ待ちになってないワープとか。順番に実行するわけではないらしい）

つまり、並列処理の中身について、1ワープ内の32スレッドと複数SMの実行は同時処理だが、同一SM内の複数ワープは同時に処理されない。
SM 内で複数ワープを保持できるので、メモリ待ちや分岐があっても次の実行可能ワープに切替される。すなわち、同じスレッド数でも、各 SM 内で保持できるワープ数が多い方が効率が良い。
（レジスタの枯渇等のハードウェア側の要因で頭打ちになる。）

サンプル : 0_Introduction/matrixMul で演算能力を確かめる

cd cuda-samples/Samples/0_Introduction/matrixMul
mkdir build
cd build
cmake .. -DCMAKE_CUDA_ARCHITECTURES=75
make
./matrixMul

実行結果　↓　

>>>
[Matrix Multiply Using CUDA] - Starting... 
GPU Device 0: "Turing" with compute capability 7.5 
MatrixA(320,320), MatrixB(640,320) 
Computing result using CUDA Kernel... 
done 
Performance= 71.64 GFlop/s, Time= 1.830 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block 
Checking computed result for correctness: Result = PASS 
NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.

理論値との比較

T400 4GB 製品情報 より理論値は 1.09TFLOP とのことなので、

71.64 GFlop/s = 0.07164 TFLOPS
0.07164 ÷ 1.09 ≈ 6.57%

約6.6%しかGPUを使えていないということが分かります。

サンプル : 4_CUDA_Libraries/matrixMulCUBLAS

数値計算ライブラリ cuBLAS を使うことでより高速に計算することができます。
cuBLAS は CUDA Toolkit に含まれている標準ライブラリです。

cd cuda-samples/Samples/4_CUDA_Libraries/matrixMulCUBLAS
mkdir build
cd build
cmake .. -DCMAKE_CUDA_ARCHITECTURES=75
make
./matrixMulCUBLAS

実行結果　↓　

>>>
[Matrix Multiply CUBLAS] - Starting...
GPU Device 0: "Turing" with compute capability 7.5

GPU Device 0: "NVIDIA T400 4GB" with compute capability 7.5

MatrixA(640,480), MatrixB(480,320), MatrixC(640,320)
Computing result using CUBLAS...done.
Performance= 243.99 GFlop/s, Time= 0.806 msec, Size= 196608000 Ops
Computing result using host CPU...done.
Comparing CUBLAS Matrix Multiply with CPU results: PASS

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

理論値との比較

243.99 GFlop/s = 0.24399 TFLOPS
0.24399 ÷ 1.09 ≈ 22.4%

6.6% と比べるとかなり効率が上がりました。

cuDNN

cuDNN とはディープ・ニューラル・ネットワーク用のライブラリです。
以下のサイトより clone して makefile の変更、OpenCV のインストールまでは問題なく行えました。
ですが wsl 環境の問題なのか、そこから先に進めませんでした。（要検討）
https://github.com/pjreddie/darknet

git clone https://github.com/pjreddie/darknet.git

# OpenCV のインストール
sudo dnf install -y opencv opencv-devel