[C++][CUDA][UE4]CUDA関数をUnreal Engine 4で用いる

ゲームエンジンの強力なグラフィックス能力は、ゲームや映像の範疇を越え、「空間」に関するあらゆるデータの描画に用いることが出来ます。医療もその1つです。

ただ、ゲームエンジンは描画以外のロジック部分は基本的にはCPU上でシングルスレッドで動いており、高負荷の計算を毎フレーム行ってリアルタイムに表示する、というようなことをしたい場合にはゴリゴリ自分で実装する必要があります。

UnityであればC#で、Unreal Engine 4であればC++で実装するのが基本です。
Unreal Engine 4をWindowsマシンで用いる場合、Visual Studioに標準で入っているライブラリであれば、必要なヘッダファイルをインクルードするだけで並列化の各種ライブラリを使うことが出来ます（出来ました）。

CPUによるfor文のマルチスレッド処理であれば並列パターンライブラリ (PPL)
GPGPUによる並列処理であればC++ AMP (C++ Accelerated Massive Parallelism)

を用いれば、気軽に並列化を行うことが出来ます（出来ました）。
（但し、Unreal Engine 4特有の問題で、C++ AMPのヘッダファイルはamp.hなのですが、困ったことにUnreal Engine 4のヘッダファイルにAmp.hという同名のものがあり、amp.hを正しくインクルードするためには絶対パスを書く必要がある、というのがだいぶ落とし穴でした…。これ、Epic Gamesさんが何かしら対処されたほうが良いのでは…）

C++ AMPは特定のGPUに限定されず、また、用いているPCにGPUが積まれていない場合には自動的にCPU側に処理を行わせるようにしてくれるという点で便利なのですが、カーネル関数の呼び出しコストや、そもそもの速さについては、やはりNVIDIA専用に最適化されているCUDAには敵いません。

参考記事：
Performance comparison of CUDA, OpenCL, and C++ AMP
Poor performance with C++ AMP Algorithms Library

となるとやはりCUDAのカーネル関数を直接用いたいわけですが、CUDAで記載された.cuファイルのコンパイルにはCUDAコンパイラを使わなければならず、Unreal Engine 4の場合はまずUnrealBuildToolが働いて…など、どこをどうすれば良いのかよくわかりません…。

参考記事：
UnrealBuildTool勉強会まとめ

CUDAコンパイラから生成された中間オブジェクトファイル.objをUnreal Engine 4から作られたファイルとリンクして…など、原理的には何とかなりそうな気もしなくもないのですが、結局よくわかりませんでした…。

海外のフォーラムを調べると、CUDAコードから静的ライブラリを作って読めば良いのでは？という示唆はあるものの、具体的な方法はおそらくまだどこにも記事がありません。

参考記事：
How to use CUDA file in UE4?

FacebookのUnreal Engine 4非公式グループで聞いてみたところ、やはり同じ手法を示唆されました。

…というわけで、本腰を入れてやってみたところ出来ました！

実行環境は、Windows 10, Unreal Engine 4.16.3, Visual Studio 2017です。

◆CUDAをインストール
Visual StudioをインストールしたあとでCUDAをインストールします。
一昔前はCUDA環境の設定だけで一苦労だったようなのですが、いまは下記からダウンロードすれば全部勝手にやってくれます。
What’s New in CUDA

◆CUDAの静的ライブラリ（.lib）作成のための下準備
他にもやり方はいくらでもあるのだとは思いますが、下記の方法が面倒な設定もほとんどなく、一番簡単でした。

1. CUDAプロジェクトの作成

CUDA 9.0 Runtimeを選びます。プロジェクト名は適当に決めて下さい。何でも良いでしょう。ここではcuda_lib_testとしました。

2. 静的ライブラリ生成用の設定

DebugモードからReleaseモードに切り替えます。Debugモードのまま静的ライブラリを作ってしまうと、Unreal Engine 4で使うことが出来ません！

ConfigurationがReleaseになっていることを確認して、Configuration TypeをApplication (.exe)からStatic library (.lib)に切り替えます。

これで設定は完了です。

◆CUDA関数の実装
デフォルトで、kernel.cuというCUDAのサンプル関数が自動的に作られますので、最低限の編集をします。
CUDAが独自に用意している型をUnreal Engine 4上からも問題なく使えるか確かめたかったので、addWithCudaの他にint4型を用いるaddWithCuda2も作ってみました。
当然ですが、ヘッダファイルに切り出す作業も必要です。以下のようにしました。

// Copyright SCIEMENT, Inc.
// by Hirofumi Seo, M.D., CEO & President

#include "cuda_lib_test.h"


__global__ void addKernel(int *c, const int *a, const int *b)
{
  int i = threadIdx.x;
  c[i] = a[i] + b[i];
}

__global__ void addKernel2(int4 *c, const int4 *a, const int4 *b)
{
  int i = threadIdx.x;
  c[i].x = a[i].x + b[i].x;
  c[i].y = a[i].y + b[i].y;
  c[i].z = a[i].z + b[i].z;
  c[i].w = a[i].w + b[i].w;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size, std::string* error_message)
{
  int *dev_a = 0;
  int *dev_b = 0;
  int *dev_c = 0;
  cudaError_t cuda_status;

  // Choose which GPU to run on, change this on a multi-GPU system.
  cuda_status = cudaSetDevice(0);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?";
    goto Error;
  }

  // Allocate GPU buffers for three vectors (two input, one output)    .
  cuda_status = cudaMalloc((void**)&dev_c, size * sizeof(int));
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMalloc failed!";
    goto Error;
  }

  cuda_status = cudaMalloc((void**)&dev_a, size * sizeof(int));
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMalloc failed!";
    goto Error;
  }

  cuda_status = cudaMalloc((void**)&dev_b, size * sizeof(int));
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMalloc failed!";
    goto Error;
  }

  // Copy input vectors from host memory to GPU buffers.
  cuda_status = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMemcpy failed!";
    goto Error;
  }

  cuda_status = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMemcpy failed!";
    goto Error;
  }

  // Launch a kernel on the GPU with one thread for each element.
  addKernel <<<1, size>>> (dev_c, dev_a, dev_b);

  // Check for any errors launching the kernel
  cuda_status = cudaGetLastError();
  if (cuda_status != cudaSuccess) {
    *error_message = "addKernel launch failed: " + std::string(cudaGetErrorString(cuda_status));
    goto Error;
  }

  // cudaDeviceSynchronize waits for the kernel to finish, and returns
  // any errors encountered during the launch.
  cuda_status = cudaDeviceSynchronize();
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaDeviceSynchronize returned error code " + std::to_string(cuda_status) +" after launching addKernel!";
    goto Error;
  }

  // Copy output vector from GPU buffer to host memory.
  cuda_status = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMemcpy failed!";
    goto Error;
  }

Error:
  cudaFree(dev_c);
  cudaFree(dev_a);
  cudaFree(dev_b);

  return cuda_status;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda2(int4 *c, const int4 *a, const int4 *b, std::string* error_message)
{
  int4 *dev_a = 0;
  int4 *dev_b = 0;
  int4 *dev_c = 0;
  const unsigned int size = 1;
  cudaError_t cuda_status;

  // Choose which GPU to run on, change this on a multi-GPU system.
  cuda_status = cudaSetDevice(0);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?";
    goto Error;
  }

  // Allocate GPU buffers for three vectors (two input, one output)    .
  cuda_status = cudaMalloc((void**)&dev_c, size * sizeof(int4));
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMalloc failed!";
    goto Error;
  }

  cuda_status = cudaMalloc((void**)&dev_a, size * sizeof(int4));
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMalloc failed!";
    goto Error;
  }

  cuda_status = cudaMalloc((void**)&dev_b, size * sizeof(int4));
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMalloc failed!";
    goto Error;
  }

  // Copy input vectors from host memory to GPU buffers.
  cuda_status = cudaMemcpy(dev_a, a, size * sizeof(int4), cudaMemcpyHostToDevice);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMemcpy failed!";
    goto Error;
  }

  cuda_status = cudaMemcpy(dev_b, b, size * sizeof(int4), cudaMemcpyHostToDevice);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMemcpy failed!";
    goto Error;
  }

  // Launch a kernel on the GPU with one thread for each element.
  addKernel2 <<<1, size>>> (dev_c, dev_a, dev_b);

  // Check for any errors launching the kernel
  cuda_status = cudaGetLastError();
  if (cuda_status != cudaSuccess) {
    *error_message = "addKernel launch failed: " + std::string(cudaGetErrorString(cuda_status));
    goto Error;
  }

  // cudaDeviceSynchronize waits for the kernel to finish, and returns
  // any errors encountered during the launch.
  cuda_status = cudaDeviceSynchronize();
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaDeviceSynchronize returned error code " + std::to_string(cuda_status) + " after launching addKernel!";
    goto Error;
  }

  // Copy output vector from GPU buffer to host memory.
  cuda_status = cudaMemcpy(c, dev_c, size * sizeof(int4), cudaMemcpyDeviceToHost);
  if (cuda_status != cudaSuccess) {
    *error_message = "cudaMemcpy failed!";
    goto Error;
  }

Error:
  cudaFree(dev_c);
  cudaFree(dev_a);
  cudaFree(dev_b);

  return cuda_status;
}

100

101

102

103

104

105

106

107

108

109

110

111

112

113

114

115

116

117

118

119

120

121

122

123

124

125

126

127

128

129

130

131

132

133

134

135

136

137

138

139

140

141

142

143

144

145

146

147

148

149

150

151

152

153

154

155

156

157

158

159

160

161

162

163

164

165

166

167

168

169

170

171

172

173

174

175

176

177

178

179

180

181

// Copyright SCIEMENT, Inc.

// by Hirofumi Seo, M.D., CEO & President

#include "cuda_lib_test.h"

__global__ void addKernel(int *c, const int *a, const int *b)

{

int i = threadIdx.x;

c[i] = a[i] + b[i];

}

__global__ void addKernel2(int4 *c, const int4 *a, const int4 *b)

{

int i = threadIdx.x;

c[i].x = a[i].x + b[i].x;

c[i].y = a[i].y + b[i].y;

c[i].z = a[i].z + b[i].z;

c[i].w = a[i].w + b[i].w;

}

// Helper function for using CUDA to add vectors in parallel.

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size, std::string* error_message)

{

int *dev_a = 0;

int *dev_b = 0;

int *dev_c = 0;

cudaError_t cuda_status;

// Choose which GPU to run on, change this on a multi-GPU system.

cuda_status = cudaSetDevice(0);

if (cuda_status != cudaSuccess) {

*error_message = "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?";

goto Error;

}

// Allocate GPU buffers for three vectors (two input, one output) .

cuda_status = cudaMalloc((void**)&dev_c, size * sizeof(int));

if (cuda_status != cudaSuccess) {

*error_message = "cudaMalloc failed!";

goto Error;

}

cuda_status = cudaMalloc((void**)&dev_a, size * sizeof(int));

if (cuda_status != cudaSuccess) {

*error_message = "cudaMalloc failed!";

goto Error;

}

cuda_status = cudaMalloc((void**)&dev_b, size * sizeof(int));

if (cuda_status != cudaSuccess) {

*error_message = "cudaMalloc failed!";

goto Error;

}

// Copy input vectors from host memory to GPU buffers.

cuda_status = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);

if (cuda_status != cudaSuccess) {

*error_message = "cudaMemcpy failed!";

goto Error;

}

cuda_status = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

if (cuda_status != cudaSuccess) {

*error_message = "cudaMemcpy failed!";

goto Error;

}

// Launch a kernel on the GPU with one thread for each element.

addKernel <<<1, size>>> (dev_c, dev_a, dev_b);

// Check for any errors launching the kernel

cuda_status = cudaGetLastError();

if (cuda_status != cudaSuccess) {

*error_message = "addKernel launch failed: " + std::string(cudaGetErrorString(cuda_status));

goto Error;

}

// cudaDeviceSynchronize waits for the kernel to finish, and returns

// any errors encountered during the launch.

cuda_status = cudaDeviceSynchronize();

if (cuda_status != cudaSuccess) {

*error_message = "cudaDeviceSynchronize returned error code " + std::to_string(cuda_status) +" after launching addKernel!";

goto Error;

}

// Copy output vector from GPU buffer to host memory.

cuda_status = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

if (cuda_status != cudaSuccess) {

*error_message = "cudaMemcpy failed!";

goto Error;

}

Error:

cudaFree(dev_c);

cudaFree(dev_a);

cudaFree(dev_b);

return cuda_status;

}

// Helper function for using CUDA to add vectors in parallel.

cudaError_t addWithCuda2(int4 *c, const int4 *a, const int4 *b, std::string* error_message)

{

int4 *dev_a = 0;

int4 *dev_b = 0;

int4 *dev_c = 0;

const unsigned int size = 1;

cudaError_t cuda_status;

// Choose which GPU to run on, change this on a multi-GPU system.

cuda_status = cudaSetDevice(0);

if (cuda_status != cudaSuccess) {

*error_message = "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?";

goto Error;

}

// Allocate GPU buffers for three vectors (two input, one output) .

cuda_status = cudaMalloc((void**)&dev_c, size * sizeof(int4));

if (cuda_status != cudaSuccess) {

*error_message = "cudaMalloc failed!";

goto Error;

}

cuda_status = cudaMalloc((void**)&dev_a, size * sizeof(int4));

if (cuda_status != cudaSuccess) {

*error_message = "cudaMalloc failed!";

goto Error;

}

cuda_status = cudaMalloc((void**)&dev_b, size * sizeof(int4));

if (cuda_status != cudaSuccess) {

*error_message = "cudaMalloc failed!";

goto Error;

}

// Copy input vectors from host memory to GPU buffers.

cuda_status = cudaMemcpy(dev_a, a, size * sizeof(int4), cudaMemcpyHostToDevice);

if (cuda_status != cudaSuccess) {

*error_message = "cudaMemcpy failed!";

goto Error;

}

cuda_status = cudaMemcpy(dev_b, b, size * sizeof(int4), cudaMemcpyHostToDevice);

if (cuda_status != cudaSuccess) {

*error_message = "cudaMemcpy failed!";

goto Error;

}

// Launch a kernel on the GPU with one thread for each element.

addKernel2 <<<1, size>>> (dev_c, dev_a, dev_b);

// Check for any errors launching the kernel

cuda_status = cudaGetLastError();

if (cuda_status != cudaSuccess) {

*error_message = "addKernel launch failed: " + std::string(cudaGetErrorString(cuda_status));

goto Error;

}

// cudaDeviceSynchronize waits for the kernel to finish, and returns

// any errors encountered during the launch.

cuda_status = cudaDeviceSynchronize();

if (cuda_status != cudaSuccess) {

*error_message = "cudaDeviceSynchronize returned error code " + std::to_string(cuda_status) + " after launching addKernel!";

goto Error;

}

// Copy output vector from GPU buffer to host memory.

cuda_status = cudaMemcpy(c, dev_c, size * sizeof(int4), cudaMemcpyDeviceToHost);

if (cuda_status != cudaSuccess) {

*error_message = "cudaMemcpy failed!";

goto Error;

}

Error:

cudaFree(dev_c);

cudaFree(dev_a);

cudaFree(dev_b);

return cuda_status;

}

// Copyright SCIEMENT, Inc.
// by Hirofumi Seo, M.D., CEO & President

#pragma once

#include <string>
#include "cuda_runtime.h"
#include "vector_types.h"
#include "vector_functions.h"
#include "device_launch_parameters.h"

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size, std::string* error_message);
cudaError_t addWithCuda2(int4 *c, const int4 *a, const int4 *b, std::string* error_message);

// Copyright SCIEMENT, Inc.

// by Hirofumi Seo, M.D., CEO & President

#pragma once

#include <string>

#include "cuda_runtime.h"

#include "vector_types.h"

#include "vector_functions.h"

#include "device_launch_parameters.h"

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size, std::string* error_message);

cudaError_t addWithCuda2(int4 *c, const int4 *a, const int4 *b, std::string* error_message);

Unreal Engine 4でのfprintfは良くわからないので（知識がないだけですが…）、UE_LOGで表示出来るようにstd::stringを作ってエラー表示を
fprintf(stderr, "xxxxx");
から
*error_message = "xxxxx";
に変更しました。
あとは基本的にはサンプルコードそのままです。

◆CUDAの静的ライブラリ（.lib）の作成

ビルドするだけです。エラーがなければ、これで.libが作られます。

ProjectPath/x64/Releaseフォルダに.libが作られています。

◆Unreal Engine 4プロジェクトの作成
何でも良いので作りましょう。CUDATestというC++ Projectを作りました。
そしてコードを書くために適当にC++ Classを作ります。CUDA_testというActor継承クラスを作りました。

そして、ヘッダファイルと先ほど作った.libファイルを該当のUnreal Engine 4プロジェクトのフォルダ下層にコピーします。もちろんコピーしなくても指定すれば使えますが、コピーしておいたほうが何かと手間が省けます。
ここではCUDALibフォルダを作り、さらにincludeフォルダ、libフォルダを作り、ヘッダファイルと.libファイルをそれぞれコピーしました。

◆Visual Studioの設定
一つだけやるとすれば、プロジェクトのPropertiesでVC++ DirectoriesのInclude Directoriesに先ほど作ったincludeフォルダを追加することでしょうか。こうするとIntelliSenseが働いてくれます。

◆Build.csの設定
ここが一番大事です。

// Copyright SCIEMENT, Inc.
// by Hirofumi Seo, M.D., CEO & President

using UnrealBuildTool;
using System.IO;

public class CUDATest : ModuleRules
{
    private string poject_root_path
    {
        get { return Path.Combine(ModuleDirectory, "../.."); }
    }

    public CUDATest(ReadOnlyTargetRules Target) : base(Target)
    {
        PCHUsage = PCHUsageMode.UseExplicitOrSharedPCHs;

        PublicDependencyModuleNames.AddRange(new string[] { "Core", "CoreUObject", "Engine", "InputCore" });

        PrivateDependencyModuleNames.AddRange(new string[] { });


        string custom_cuda_lib_include = "CUDALib/include";
        string custom_cuda_lib_lib = "CUDALib/lib";

        PublicIncludePaths.Add(Path.Combine(poject_root_path, custom_cuda_lib_include));
        PublicAdditionalLibraries.Add(Path.Combine(poject_root_path, custom_cuda_lib_lib, "cuda_lib_test.lib"));

        string cuda_path = "C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0";
        string cuda_include = "include";
        string cuda_lib = "lib/x64";

        PublicIncludePaths.Add(Path.Combine(cuda_path, cuda_include));

        //PublicAdditionalLibraries.Add(Path.Combine(cuda_path, cuda_lib, "cudart.lib"));
        PublicAdditionalLibraries.Add(Path.Combine(cuda_path, cuda_lib, "cudart_static.lib"));

        // Uncomment if you are using Slate UI
        // PrivateDependencyModuleNames.AddRange(new string[] { "Slate", "SlateCore" });

        // Uncomment if you are using online features
        // PrivateDependencyModuleNames.Add("OnlineSubsystem");

        // To include OnlineSubsystemSteam, add it to the plugins section in your uproject file with the Enabled attribute set to true
    }
}

// Copyright SCIEMENT, Inc.

// by Hirofumi Seo, M.D., CEO & President

using UnrealBuildTool;

using System.IO;

public class CUDATest : ModuleRules

{

private string poject_root_path

{

get { return Path.Combine(ModuleDirectory, "../.."); }

}

public CUDATest(ReadOnlyTargetRules Target) : base(Target)

{

PCHUsage = PCHUsageMode.UseExplicitOrSharedPCHs;

PublicDependencyModuleNames.AddRange(new string[] { "Core", "CoreUObject", "Engine", "InputCore" });

PrivateDependencyModuleNames.AddRange(new string[] { });

string custom_cuda_lib_include = "CUDALib/include";

string custom_cuda_lib_lib = "CUDALib/lib";

PublicIncludePaths.Add(Path.Combine(poject_root_path, custom_cuda_lib_include));

PublicAdditionalLibraries.Add(Path.Combine(poject_root_path, custom_cuda_lib_lib, "cuda_lib_test.lib"));

string cuda_path = "C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0";

string cuda_include = "include";

string cuda_lib = "lib/x64";

PublicIncludePaths.Add(Path.Combine(cuda_path, cuda_include));

//PublicAdditionalLibraries.Add(Path.Combine(cuda_path, cuda_lib, "cudart.lib"));

PublicAdditionalLibraries.Add(Path.Combine(cuda_path, cuda_lib, "cudart_static.lib"));

// Uncomment if you are using Slate UI

// PrivateDependencyModuleNames.AddRange(new string[] { "Slate", "SlateCore" });

// Uncomment if you are using online features

// PrivateDependencyModuleNames.Add("OnlineSubsystem");

// To include OnlineSubsystemSteam, add it to the plugins section in your uproject file with the Enabled attribute set to true

}

27行目までが自分で作った.libとヘッダファイルを指定しているところで、29行目からがCUDA用の設定です。
CUDA 9.0の場合、標準的なインストールでは
C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0
以下にヘッダファイルや.libファイルが置かれているはずです。
CUDA 9.0の.libは、cudart_static.libでも、コメントアウトしているcudart.libでもどちらでも大丈夫です。両方指定するとどうなるのかはよくわかりません（未検証）。
この2つの違いは例えば以下に書かれています。
cudart.lib vs. cudart_static.lib – NVIDIA Developer Forums
cudart_static.libは、その名の通り静的ライブラリで、.exeにした際に.dllが不要になります。

CUDA_test Actorは以下のようにしてみました。検証用なので実装はヘッダファイルに書きました。

// Copyright SCIEMENT, Inc.
// by Hirofumi Seo, M.D., CEO & President

#pragma once

#include "CoreMinimal.h"
#include "cuda_lib_test.h"
#include "GameFramework/Actor.h"
#include "cuda_runtime.h"
#include "CUDA_test.generated.h"

UCLASS()
class CUDATEST_API ACUDA_test : public AActor
{
  GENERATED_BODY()

public:
  // Sets default values for this actor's properties
  ACUDA_test();

  UFUNCTION(BlueprintCallable, Category = "CUDATest")
  bool SimpleCUDATest() {
    // ----- addWithCuda test -----
    const int arraySize = 5;
    const int a[arraySize] = { 1, 2, 3, 4, 5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };
    std::string error_message;

    // Add vectors in parallel.
    cudaError_t cuda_status = addWithCuda(c, a, b, arraySize, &error_message);
    if (cuda_status != cudaSuccess) {
      UE_LOG(LogTemp, Warning, TEXT("addWithCuda failed!\n"));
      UE_LOG(LogTemp, Warning, TEXT("%s"), *FString(error_message.c_str()));
      return false;
    }
    UE_LOG(LogTemp, Warning, TEXT("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}"), c[0], c[1], c[2], c[3], c[4]);

    // ----- addWithCuda2 test -----
    const int4 a_int4 = make_int4(1, 2, 3, 4);
    const int4 b_int4 = make_int4(10, 20, 30, 40);
    int4 c_int4;

    // Add vectors in parallel.
    cuda_status = addWithCuda2(&c_int4, &a_int4, &b_int4, &error_message);
    if (cuda_status != cudaSuccess) {
      UE_LOG(LogTemp, Warning, TEXT("addWithCuda failed!\n"));
      UE_LOG(LogTemp, Warning, TEXT("%s"), *FString(error_message.c_str()));
      return false;
    }
    UE_LOG(LogTemp, Warning, TEXT("{1,2,3,4} + {10,20,30,40} = {%d,%d,%d,%d}"), c_int4.x, c_int4.y, c_int4.z, c_int4.w);

    return true;
  }

protected:
  // Called when the game starts or when spawned
  virtual void BeginPlay() override;

public:
  // Called every frame
  virtual void Tick(float DeltaTime) override;

};

// Copyright SCIEMENT, Inc.

// by Hirofumi Seo, M.D., CEO & President

#pragma once

#include "CoreMinimal.h"

#include "cuda_lib_test.h"

#include "GameFramework/Actor.h"

#include "cuda_runtime.h"

#include "CUDA_test.generated.h"

UCLASS()

class CUDATEST_API ACUDA_test : public AActor

{

GENERATED_BODY()

public:

// Sets default values for this actor's properties

ACUDA_test();

UFUNCTION(BlueprintCallable, Category = "CUDATest")

bool SimpleCUDATest() {

// ----- addWithCuda test -----

const int arraySize = 5;

const int a[arraySize] = { 1, 2, 3, 4, 5 };

const int b[arraySize] = { 10, 20, 30, 40, 50 };

int c[arraySize] = { 0 };

std::string error_message;

// Add vectors in parallel.

cudaError_t cuda_status = addWithCuda(c, a, b, arraySize, &error_message);

if (cuda_status != cudaSuccess) {

UE_LOG(LogTemp, Warning, TEXT("addWithCuda failed!\n"));

UE_LOG(LogTemp, Warning, TEXT("%s"), *FString(error_message.c_str()));

return false;

}

UE_LOG(LogTemp, Warning, TEXT("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}"), c[0], c[1], c[2], c[3], c[4]);

// ----- addWithCuda2 test -----

const int4 a_int4 = make_int4(1, 2, 3, 4);

const int4 b_int4 = make_int4(10, 20, 30, 40);

int4 c_int4;

// Add vectors in parallel.

cuda_status = addWithCuda2(&c_int4, &a_int4, &b_int4, &error_message);

if (cuda_status != cudaSuccess) {

UE_LOG(LogTemp, Warning, TEXT("addWithCuda failed!\n"));

UE_LOG(LogTemp, Warning, TEXT("%s"), *FString(error_message.c_str()));

return false;

}

UE_LOG(LogTemp, Warning, TEXT("{1,2,3,4} + {10,20,30,40} = {%d,%d,%d,%d}"), c_int4.x, c_int4.y, c_int4.z, c_int4.w);

return true;

}

protected:

// Called when the game starts or when spawned

virtual void BeginPlay() override;

public:

// Called every frame

virtual void Tick(float DeltaTime) override;

};

これでビルドが通ります！！！

◆Unreal Editor内での設定
特に特別なことは必要ありません。
ここではCUDA_testクラスからBlueprintクラスを作り、クラス関数で先ほど作ったSimpleCUDATest()を呼び出すようにしました。
これをシーン内に配置します。

実行します。Window -> Developer Tools -> Output LogでOutput Logを出しておきましょう。Play!

CUDAのカーネル関数が呼び出されて、足し算が正しく計算出来ていることがわかります！
やったねっ！！！

…というわけで、正直これではGPUの性能を何一つ活かしていないわけですが、しかしこれでUnreal Engine 4でCUDAを操ることが出来ます！
大手の会社や有償サポートコミュニティ内では当たり前なのかもしれませんが、公になっている記事はおそらくこれが世界初だと思います。

大規模計算のリアルタイム可視化などにUnreal Engine 4を使うのも良いのではないでしょうか？
（そんな案件がありましたらぜひ！サイアメントのContactページからお気軽にお問い合わせください。）

※本記事内容は、国立研究開発法人日本医療研究開発機構（AMED）の平成29年度「未来医療を実現する医療機器・システム研究開発事業『術中の迅速な判断・決定を支援するための診断支援機器・システム開発』」採択課題である「術前と術中をつなぐスマート手術ガイドソフトウェアの開発」（代表機関名：東京大学、研究開発代表者名：齊藤延人）に、東京大学大学院情報理工学系研究科の学術支援専門職員として参画している瀬尾拡史が、研究開発として行っているものやその成果を一部含んでいます。

Post Views: 99,219

3 Comments

Saurabh Sharma より:

2018年2月3日 4:56 AM

Hey nice article , the only thing i am a bit confused is what about launching the kernel like in .cu file you have a kernel named addKernel2, what if i want to increase the number of threads per block like say 256 threads . For this do i have to create a cu file again and make another .lib file then i have to repeat the whole process to make it wotk with unreal? Or is there a way i could directly call __global__ function in ue4?

返信
1. Ross Harrison より:
  
  2019年1月4日 2:28 PM
  
  My guess is yes, but it’s more difficult.
  
  You would have to construct your API such that you could:
  1. Read information about the coda cores from your library
  and
  2. Control the thread allocation parameters from cuda.
  
  It would be best to handle this within the cuda library, so the API can be simpler.
  
  返信
David N より:

2018年9月7日 5:41 AM

Thank you so much for the article, it worked like a charm! I’ve had to google translate it to English, but all the more fun 🙂

返信

サイアメント技術メモ

[C++][CUDA][UE4]CUDA関数をUnreal Engine 4で用いる

3 Comments

コメントを残すコメントをキャンセル

3 Comments

コメントを残す コメントをキャンセル

コメントを残すコメントをキャンセル