はじめに
CUDAはC/C++のみをサポートしますが，C++/CLIを使うことでC＃でもCUDAによる計算を簡単に実装できます．C++/CLIとはC++にC#のマネージドな機能を追加した言語で，C++やCUDAで記述したコードを，C#で呼び出せる形式のDLLとしてビルドできます．そのため下図のようにソリューション内でプロジェクトを2つに分けて, プロジェクト1でCUDAを使った関数をC++/CLIで記述し，プロジェクト2（C#側）で呼び出すことでC#でもGPUによる並列計算処理を実行できます．
そこで本稿では，C++/CLIを活用して，C#においてCUDAでベクトル和を計算するサンプルプログラムを作成します．
ソースコード：https://kkaneko-lab.com/?p=344
1. 事前準備
①Visual studio communityをインストールします．
②CUDAのインストールします．今回はCUDA Toolkit 12.3で実装します．https://developer.nvidia.com/cuda-toolkit
③ワークロードで「C++によるデスクトップ開発」をインストールします．
加えて，C++/CLIを使用するために，Visual studio Installerの個別コンポーネントでビルドツールをインストールします．CUDA 12.3では MSVC v14.3Xまでしかサポートされていないの下図の「v143 ビルドツールのC++/CLIサポート（14.39-17.9）」を追加しました．
2. プロジェクト1（C++/CLI）の作成
最初に，CUDA関数を格納するDLLを作成するためのC++/CLIプロジェクトを作成します．下図のように，CUDA 12.3 Runtimeを選択して，まずはCUDAに対応したC++のプロジェクトを作成します．今回は，”CUDACLITest”という名前で作成しました．このC++(CUDA)のプロジェクトの設定を変更することで，C++/CLIの機能を使えるようにします．
DLLを作成するためのプロジェクトなので，プロジェクトのプロパティで「構成の種類」をDLLに変更します．
次にC++/CLIプロパティを以下のように設定することで，.Net 8.0をターゲットとしたDLLの作成が可能となります．
念のため，MSVCツールセットのバージョンを確認します．上記の通り，CUDA 12.3 ではMSVCはv14.3Xまでしかサポートされていないので，もしこれが14.4Xとなっている場合は，以下のように14.3Xに変更しないとビルドが通らない可能性があります．
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
#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はヘッダーファイルで，以下の通りです．
#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);
#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<float>^ vec_0, array<float>^ vec_1)を記述します．今回は，ネームスペースをCUDA，クラス名をCalculatorとしました．クラスの定義において，public ref class となっておりますが，refを付けることでC#で呼び出し可能なクラスとして定義できます．
以下のようにCalculator.cppを記述しました．この関数の引数の型は，array<float>ですがこれは，C#におけるfloat[]と同義です．このようにC++/CLIではC#のマネージドな機能を使用できます．マネージドな変数には頭に^が付きます．
関数 “Calculator::Add” では，まずC#のマネージド配列 array を C++ のアンマネージド配列 float* に pin_ptr を使って変換します．次はデバイス変数のメモリ確保とデータ転送を行っていますが，個々の処理は通常のCUDAと同様です．その後，カーネル関数を呼び出してGPU上で計算を行いますが，C++/CLI では通常の CUDA と異なり ”cudaLaunchKernel” によりカーネルを実行します．最後にデバイスでの計算結果 d_result をホスト変数 h_result に転送します．そしてマネージド配列に変換して計算結果を返します．
以上がC++/CLIによる関数の実装になります．このプロジェクトをビルドすることでC#でも呼び出し可能なDLLを作成できます．
4. プロジェクト2(C#)を作成
以上のように実装したCalculatorクラスを使用するC#のプロジェクトを作成します．以下のように今回はコンソールアプリを選択しました．
.Netのバージョンは，C++/CLIの設定に合わせて 8.0 と設定します．
今回は「TestConsole」という名前でプロジェクトを作成しました．作成後は以下のように参照の追加から，事前に作成したC++/CLIのプロジェクトを参照追加します．
Program.csには以下のように，自作したCalculatorクラスを使ってベクトル和をGPUで計算するコードを記述しました．正しく依存関係が設定できていれば， ”using CUDA”のようにC++/CLIのプロジェクトで定義した名前空間が認識されます．
このコードでは，長さ1000のベクトルを用意して，その和を計算します．その後，計算結果のベクトルの要素を順番にコンソール表示しています．
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ベクトルの和を計算します．
プログラムを実行した結果，コンソールには上図と同様の結果が表示されていることから，GPUを使って正しく足し算が実行できていることがわかります．
補足
C#のプロジェクト側のプロパティで，下のようにネイティブコードのデバッグを有効にすると，C++/CLIプロジェクトに設定したブレークポイントで停止することができます．
