【CUDA + .Net】C++/CLIを使ってC#でもGPUコンピューティング

# はじめに
　CUDAはC/C++のみをサポートしますが，C++/CLIを使うことでC＃でもCUDAによる計算を簡単に実装できます．C++/CLIとはC++にC#のマネージドな機能を追加した言語で，C++やCUDAで記述したコードを，C#で呼び出せる形式のDLLとしてビルドできます．そのため下図のようにソリューション内でプロジェクトを2つに分けて, プロジェクト1でCUDAを使った関数をC++/CLIで記述し，プロジェクト2（C#側）で呼び出すことで**C#でもGPUによる並列計算処理を実行できます**．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/a3e3e95e-7fb0-4c58-ed6f-66ff11e4cd96.png)

　そこで本稿では，<font color=red>**C++/CLIを活用して，C#においてCUDAでベクトル和を計算**</font>するサンプルプログラムを作成します．

 
**<font color=red>ソースコード：https://kkaneko-lab.com/?p=344 </font>**


![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/89d3a151-fc45-fcc6-17e2-a807f6c1895d.png)
 

 


# 1. 事前準備
①Visual studio communityをインストールします．
②CUDAのインストールします．今回はCUDA Toolkit 12.3で実装します．https://developer.nvidia.com/cuda-toolkit
③ワークロードで「C++によるデスクトップ開発」をインストールします．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/1c3b1d5f-5bda-71ec-f78b-e9175e31a047.png)
　加えて，C++/CLIを使用するために，Visual studio Installerの個別コンポーネントでビルドツールをインストールします．CUDA 12.3では MSVC v14.3Xまでしかサポートされていないの下図の「v143 ビルドツールのC++/CLIサポート（14.39-17.9）」を追加しました．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/0cfd4478-a36f-86df-d9c4-d96ab955dd37.png)

# 2. プロジェクト1（C++/CLI）の作成
　最初に，CUDA関数を格納するDLLを作成するためのC++/CLIプロジェクトを作成します．下図のように，CUDA 12.3 Runtimeを選択して，まずはCUDAに対応したC++のプロジェクトを作成します．今回は，”CUDACLITest”という名前で作成しました．このC++(CUDA)のプロジェクトの設定を変更することで，C++/CLIの機能を使えるようにします．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/a69bef48-5740-81aa-f8c9-3f9814b921a0.png)

　DLLを作成するためのプロジェクトなので，プロジェクトのプロパティで「構成の種類」をDLLに変更します．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/1115a9d0-ecac-542f-4adc-c9a85a9ccf56.png)

　次にC++/CLIプロパティを以下のように設定することで，.Net 8.0をターゲットとしたDLLの作成が可能となります．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/c848db22-109c-098d-aada-2a756b28c1b6.png)


　念のため，MSVCツールセットのバージョンを確認します．上記の通り，CUDA 12.3 ではMSVCはv14.3Xまでしかサポートされていないので，もしこれが14.4Xとなっている場合は，以下のように14.3Xに変更しないとビルドが通らない可能性があります．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/67342421-623b-43f3-eff5-43e9a35cdec5.png)

# 3. ベクトル和を計算するCUDAコードを記述
　以下のように4つのファイルでDLLを構築します．
①Kernel.cu・・・・・ベクトル和を計算するカーネル関数を定義
②Kernel.cuh ・・・・Kernel.cuのヘッダーファイル
③Calculator.cpp・・・カーネル関数を呼び出し，計算を実行する関数を定義
④Calculator.h・・・・Calculator.cppのヘッダーファイル

Kernel.cuでは，以下のように2ベクトルの和を計算するカーネル関数を記述しました．CUDAでの並列計算の仕様は，以下のプログラミングガイドを参照してください．https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

```Kernel.cu
#pragma once
#include "Kernel.cuh"

/// <summary>
/// ベクトル和を計算するカーネル関数
/// </summary>
/// <param name="vec_0">ベクトル0</param>
/// <param name="vec_1">ベクトル1</param>
/// <param name="result">計算結果のベクトル</param>
/// <param name="length">ベクトルの長さ</param>
/// <returns></returns>
__global__ void CudaAddKernel(float* vec_0, float* vec_1, float* result, int* length) {
	//インデックス
	int index = (blockIdx.x * blockDim.x) + threadIdx.x;

	//インデックスが範囲内のとき
	if (index < *length) {
		//ベクトルの要素どうしを足し合わせる
		result[index] = vec_0[index] + vec_1[index];
	}
}
```
Kernel.cuhはヘッダーファイルで，以下の通りです．

```Kernel.cuh
#pragma
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

/// <summary>
/// ベクトル和を計算するカーネル関数
/// </summary>
/// <param name="vec_0">ベクトル0</param>
/// <param name="vec_1">ベクトル1</param>
/// <param name="result">計算結果のベクトル</param>
/// <param name="length">ベクトルの長さ</param>
/// <returns></returns>
__global__ void CudaAddKernel(float* vec_0, float* vec_1, float* result, int* length);
```

```Calculator.cpp
#pragma once
#include "Calculator.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "Kernel.cuh"

namespace CUDA {
	/// <summary>
	/// 同じ長さのベクトルの和
	/// </summary>
	/// <param name="vec_0">ベクトル0</param>
	/// <param name="vec_1">ベクトル1</param>
	/// <returns></returns>
	array<float>^ Calculator::Add(array<float>^ vec_0, array<float>^ vec_1) {
		
#pragma region ホスト変数の用意
		//pin_ptrで配列を固定し、ポインターを取得
		pin_ptr<float> vec_0_pin_ptr = &vec_0[0];
		pin_ptr<float> vec_1_pin_ptr = &vec_1[0];
		//アンマネージド配列のホスト変数を用意
		float* h_vec_0 = vec_0_pin_ptr; //ベクトル0のアンマネージド配列
		float* h_vec_1 = vec_1_pin_ptr; //ベクトル1のアンマネージド配列
		float* h_result = new float[vec_0->Length]; //計算結果のアンマネージド配列
		int h_length = vec_0->Length; //ベクトルの長さ
#pragma endregion

#pragma region デバイス変数の用意
		//ホスト変数に対応した，デバイス変数を用意
		float* d_vec_0; //ベクトル0
		float* d_vec_1; //ベクトル0
		float* d_result; //計算結果
		int* d_length; //ベクトルの長さ
		//デバイスを指定
		cudaError_t cuda_status = cudaSetDevice(0);
		//デバイス変数のメモリ確保
		cuda_status = cudaMalloc(&d_vec_0, h_length * sizeof(float));
		cuda_status = cudaMalloc(&d_vec_1, h_length * sizeof(float));
		cuda_status = cudaMalloc(&d_result, h_length * sizeof(float));
		cuda_status = cudaMalloc(&d_length, sizeof(int));
		//ホスト変数の値をデバイス変数にコピー
		cuda_status = cudaMemcpy(d_vec_0, h_vec_0, h_length * sizeof(float), cudaMemcpyHostToDevice);
		cuda_status = cudaMemcpy(d_vec_1, h_vec_1, h_length * sizeof(float), cudaMemcpyHostToDevice);
		cuda_status = cudaMemcpy(d_length, &h_length, sizeof(int), cudaMemcpyHostToDevice);
#pragma endregion

#pragma region カーネル関数の実行
		//並列計算条件
		int max_thread_num = 256; //最大スレッド数
		dim3 grid(h_length / max_thread_num + 1); //グリッドの次元
		dim3 block(max_thread_num); //ブロックの次元
		//引数
		void* args[] = { &d_vec_0, &d_vec_1, &d_result, &d_length };
		//カーネル実行
		cuda_status = cudaLaunchKernel((const void*)CudaAddKernel, grid, block, args);
		//処理待ち
		cuda_status = cudaDeviceSynchronize();
#pragma endregion

		//計算結果のデバイス変数の値を，ホストにコピー
		cuda_status = cudaMemcpy(h_result, d_result, h_length * sizeof(float), cudaMemcpyDeviceToHost);
		//アンマネージド配列を，マネージド配列に変換
		array<float>^ result_managed = gcnew array<float>(vec_0->Length);
		for (int i = 0; i < result_managed->Length; i++) result_managed[i] = h_result[i];

		//newで確保したメモリの開放
		delete[] h_result;
		//デバイスメモリの開放
		cuda_status = cudaFree(d_vec_0);
		cuda_status = cudaFree(d_vec_1);
		cuda_status = cudaFree(d_result);
		cuda_status = cudaFree(d_length);

		return result_managed;
	}
}
```

　Calculator.cppでは，上記のカーネル関数”CudaAddKernel”を呼び出してGPUで計算を行い，計算結果を返すコードCalculator::Add(array&lt;float>^ vec_0, array&lt;float>^ vec_1)を記述します．今回は，ネームスペースをCUDA，クラス名をCalculatorとしました．クラスの定義において，public ref class となっておりますが，refを付けることでC#で呼び出し可能なクラスとして定義できます．
　以下のようにCalculator.cppを記述しました．この関数の引数の型は，array&lt;float>ですがこれは，C#におけるfloat[]と同義です．このようにC++/CLIではC#のマネージドな機能を使用できます．マネージドな変数には頭に^が付きます．
　関数 “Calculator::Add” では，まずC#のマネージド配列 array<float> を C++ のアンマネージド配列 float* に pin_ptr を使って変換します．次はデバイス変数のメモリ確保とデータ転送を行っていますが，個々の処理は通常のCUDAと同様です．その後，カーネル関数を呼び出してGPU上で計算を行いますが，C++/CLI では通常の CUDA と異なり ”cudaLaunchKernel” によりカーネルを実行します．最後にデバイスでの計算結果 d_result をホスト変数 h_result に転送します．そしてマネージド配列に変換して計算結果を返します．

 　以上がC++/CLIによる関数の実装になります．このプロジェクトをビルドすることでC#でも呼び出し可能なDLLを作成できます．

  # 4. プロジェクト2(C#)を作成
  以上のように実装したCalculatorクラスを使用するC#のプロジェクトを作成します．以下のように今回はコンソールアプリを選択しました．
  ![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/f1e55046-a4e1-df9d-7331-3d8d222a769a.png)
.Netのバージョンは，C++/CLIの設定に合わせて 8.0 と設定します．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/6402608a-9540-2dd8-37f7-3618c81f21db.png)
今回は「TestConsole」という名前でプロジェクトを作成しました．作成後は以下のように参照の追加から，事前に作成したC++/CLIのプロジェクトを参照追加します．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/e02c6650-46bd-ffa1-dc7e-4dd308dbd9fb.png)
　Program.csには以下のように，自作したCalculatorクラスを使ってベクトル和をGPUで計算するコードを記述しました．正しく依存関係が設定できていれば， ”using CUDA”のようにC++/CLIのプロジェクトで定義した名前空間が認識されます．
　このコードでは，長さ1000のベクトルを用意して，その和を計算します．その後，計算結果のベクトルの要素を順番にコンソール表示しています．

```Program.cs
using CUDA;
using System.Diagnostics;

//ベクトルの長さ
int vectorLength = 1000;

//2つのベクトルを用意
float[] vector0 = new float[vectorLength];
float[] vector1 = new float[vectorLength];

//ベクトルの中身を設定
for(int i = 0; i < vectorLength; i++)
{
    vector0[i] = i;
    vector1[i] = i * 1000;
}

//GPUでベクトル和を計算
float[] sum = CUDA.Calculator.Add(vector0, vector1);

//コンソール表示
foreach(float value in sum) Console.WriteLine(value); //計算結果
```
　このプログラムでは以下のような2ベクトルの和を計算します．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/89d3a151-fc45-fcc6-17e2-a807f6c1895d.png)
プログラムを実行した結果，コンソールには上図と同様の結果が表示されていることから，GPUを使って正しく足し算が実行できていることがわかります．
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/b2893a1c-e107-8e25-3339-5fe11b185694.png)
# 補足
C#のプロジェクト側のプロパティで，下のようにネイティブコードのデバッグを有効にすると，C++/CLIプロジェクトに設定したブレークポイントで停止することができます．

![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3780116/71cb2e0f-ff01-59b7-2e15-a0412374437e.png)