Clspvを用いてOpenCLカーネルをVulkanで実行する
こんにちは。このページは レイトレ Advent Calendar 2019 の19日目の記事です。この記事では以下の事について紹介します、
- Vulkan SDK と Clspv のインストール
- Clspv で OpenCL カーネルをビルド
- ビルドした カーネルを Vulkan で実行する際の注意事項
この記事では Vulkan 仕様や API の使い方については説明しません。Vulkan Compute Shader の使い方は Clspv を使った場合でも特別な事は無いため、他の資料などを参考にしてもらえればと思います。
Clspv を使った Vulkan のサンプルコードは github にアップロードしていますので参考にしてください。
Vulkan SDK のインストール
まずは、下記のページから最新の Vulkan SDK をダウンロードします。
SDKのインストール方法は Windows、Linux、macOS で異なります。
Windows
ダウンロードしたインストーラーを起動しそのままインストールします。インストールが完了したら VULKAN_SDK という環境変数が作られ、Vulkan SDK へのパスが設定されます。
Linux
ダウンロードした tar.gz を展開します。展開したディレクトリ内に setup-env.sh というファイルがあり、これをシェルに反映させると、SDK内のツールやライブラリへのパスを自動的に設定してくれます。また、 VULKAN_SDK という環境変数がつくられて、Vulkan SDKへのパスが設定されます。
% source ${path-to-vulkan-sdk}/setup-env.sh
% echo ${VULKAN_SDK}
${path-to-vulkan-sdk}/x86_64
macOS
ダウンロードした tar.gz を展開します。Linux版と違い展開したディレクトリ内に setup-env.sh が無いため以下のようなスクリプトを用意します。
export VULKAN_SDK="$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )"/macOS
export PATH="${VULKAN_SDK}/bin:${PATH}"
export DYLD_LIBRARY_PATH="${VULKAN_SDK}/lib:${DYLD_LIBRARY_PATH}"
export VK_LAYER_PATH="${VULKAN_SDK}/etc/vulkan/explicit_layer.d"
export VK_ICD_FILENAMES="${VULKAN_SDK}/etc/vulkan/icd.d/MoltenVK_icd.json"
このスクリプトを展開したディレクトリ内に置きシェルに反映させると、SDK内のツールやライブラリへのパスを自動的に設定してくれます。
Vulkan SDK の動作確認
Vulkan SDK をインストールしたのできちんと動作するか確認します。動作確認には vkvia を用います。vkvia は簡単なVulkanアプリケーションをいくつか動かして、その時のVulkanの状態をキャプチャーします。
vkvia は ${VULKAN_SDK}/bin/ にあります。実行すると実行ディレクトリに vkvia.html が作成され、ブラウザで開くとキャプチャーしたVulkanの状態を見ることができます。例えば、 Physical Devices を開くとVulkanで使用できるハードウェアデバイスの情報を見ることができます。
Clspv のインストール
Clspvにはリリースバージョンの実行ファイルやソースコードはまだありません。github の google/clspv から master ブランチを pull してビルドします。ビルドするのに必要なツールはREADMEに書いてあるように、
- CMake
- Python
- C++ compiler
- git
になります。まずはClspvが依存する外部プロジェクトのソースコードをダウンロードします。
% python utils/fetch_sources.py
ダウンロードした後は CMake を実行して Clspv のビルドの準備をします。Windows の場合は (Visual Studio 2019 を使う場合)、
% cmake -G"Visual Studio 16 2019" -Ax64 -Thost=x64 -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=${path-to-dir} ..
${path-to-dir} にはインストール先のパスが入ります。
Linux や macOS の場合は、
% cmake -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=${path-to-install} ..
です。 CMake でジェネレーターファイルを作成したら Clspv をビルドします。
% cmake --build . --config Release --target install --parallel 8
ビルドが完了すると ${path-to-dir} に clspv という実行ファイルがインストールされます。
Clspv で OpenCL コードをビルド
Clspv は OpenCL C 1.2 のコードをビルドすることができますが、Vulkan が扱うことができない機能については Clspv 側でビルドができないように制限がかけられています。制限内容の詳細は OpenCL C Restrictions を読んで下さい。いくつかピックアップすると、
- 8や16長のベクトルは使用できない (
int8など) doubleは使用できないcbrt、pown、rint、sincosなど使用できない Math 関数があるprintfは使用できない
などがあります。ただ、これらの制限は将来 Vulkan の仕様のアップデートに合わせて緩和されるかもしれません。
では、実際に Clspv で以下の OpenCL C コードをビルドしてみます。
// vulkan_clspv_test0.cl
__kernel void test(__global const float4* inputs, __global float4* outputs)
{
const uint index = get_global_id(0);
outputs[index] = sqrt(inputs[index]);
}
コマンドは以下のようになります。
% clspv -o=vulkan_clspv_test0.spv vulkan_clspv_test0.cl
-o オプションで出力する SPIR-V ファイルのファイル名を指定できます。その他にも最適化オプション -O=3 などがあり、オプション一覧は clspv --help で見ることができます。
OpenCL C++ のビルド
Clspv は実験的ですが OpenCL C++ をサポートします。現状 OpenCL C++ の仕様を完全にサポートしているわけではく、いくつかの違いはあります。例えば、
- OpenCL C++ Standard Library は使用できない
- Mathライブラリなどは OpenCL C の組み込み関数スタイルの物を使用する
- C++14 ではなく C++17 ベース
constexpr ifやfold expressionなど C++17 の機能も使うことができる
等の違いがあります。Clspv で OpenCL C++ のコードをビルドする場合は --c++ --inline-entry-points オプションを付けてビルドします。例として、以下の OpenCL C++ コードをビルドしてみます。
// vulkan_clspv_test1.cl
namespace clspvtest {
// Forward declaration
float sum(const float x);
float sum(const float4 x);
template <typename... Types> float sum(const Types... args);
class Matrix4x4;
Matrix4x4 operator+(const Matrix4x4& lhs, const Matrix4x4& rhs);
float sum(const float x)
{
return x;
}
float sum(const float4 x)
{
const float s = x.x + x.y + x.z + x.w;
return s;
}
template <typename... Types>
float sum(const Types... args)
{
// Clspv can build C++17 features
if constexpr (0 < sizeof...(Types)) {
const auto s = (args + ...);
return sum(s);
}
else {
return 0.0f;
}
}
class Matrix4x4
{
public:
Matrix4x4() {}
Matrix4x4(const float4 r1, const float4 r2, const float4 r3, const float4 r4) :
r1_{r1}, r2_{r2}, r3_{r3}, r4_{r4}
{
}
//! Return the sum of the elements
float sum() const
{
const float s = clspvtest::sum(r1_, r2_, r3_, r4_);
return s;
}
float4 r1_ = float4{0.0f, 0.0f, 0.0f, 0.0f},
r2_ = float4{0.0f, 0.0f, 0.0f, 0.0f},
r3_ = float4{0.0f, 0.0f, 0.0f, 0.0f},
r4_ = float4{0.0f, 0.0f, 0.0f, 0.0f};
};
//! Add corresponding elements of the given two matrices
Matrix4x4 operator+(const Matrix4x4& lhs, const Matrix4x4& rhs)
{
const Matrix4x4 result{lhs.r1_ + rhs.r1_, lhs.r2_ + rhs.r2_, lhs.r3_ + rhs.r3_, lhs.r4_ + rhs.r4_};
return result;
}
} // namespace clspvtest
__kernel void testSummation(__global clspvtest::Matrix4x4* value, __global float* outputs)
{
const uint index = get_global_id(0);
if (0 < index)
return;
// Check clspvtest::Matrix4x4
static_assert(sizeof(clspvtest::Matrix4x4) == 64, "The size of clspvtest::Matrix4x4 isn't 64 bytes.");
static_assert(alignof(clspvtest::Matrix4x4) == 16, "The alignment of clspvtest::Matrix4x4 isn't 16 bytes.");
// Constant value
__constant constexpr float k = 4.0f; // constant value must be in the outermost scope of a kernel
static_assert(k == 4.0f);
outputs[0] = clspvtest::sum();
outputs[1] = clspvtest::sum(sqrt(k)); // Use OpenCL C style built-in function
outputs[2] = clspvtest::sum(1.0f, 2.0f, 3.0f);
// Matrix4x4 in registers
{
__private const clspvtest::Matrix4x4 m;
outputs[3] = m.sum();
}
{
__private const clspvtest::Matrix4x4 m{float4{k, k, k, k},
float4{k, k, k, k},
float4{k, k, k, k},
float4{k, k, k, k}};
outputs[4] = m.sum();
}
// Matrix4x4 in global memory
{
__private const clspvtest::Matrix4x4 m1{float4{k, k, k, k},
float4{k, k, k, k},
float4{k, k, k, k},
float4{k, k, k, k}};
__private const clspvtest::Matrix4x4 m2{float4{0.0f, 1.0f, 2.0f, 3.0f},
float4{4.0f, 5.0f, 6.0f, 7.0f},
float4{8.0f, 9.0f, 10.0f, 11.0f},
float4{12.0f, 13.0f, 14.0f, 15.0f}};
value[0] = m1 + m2; // Initialize global Matrix4x4
outputs[5] = value->sum();
}
// Matrix4x4 in local memory
__local clspvtest::Matrix4x4 storage[1]; // local variable must be in the outermost scope of a kernel
{
__private const clspvtest::Matrix4x4 m = value[0];
storage[0] = m;
outputs[6] = storage->sum();
}
}
以下のコマンドでビルドします。
% clspv --c++ --inline-entry-points -o=vulkan_clspv_test1.spv vulkan_clspv_test1.cl
ビルドが成功して SPIR-V ファイルが生成されます。実行結果も以下のように、
output[0] = 0
output[1] = 2
output[2] = 6
output[3] = 0
output[4] = 64
output[5] = 184
output[6] = 184
正しく実行されています。( github のサンプルコードをビルドすると実行できます。)
Vulkan で OpenCL カーネルを実行する
Clspv で生成した SPIR-V コードは Vulkan Compute Shader で実行できます。SPIR-V コードの実行方法は通常の方法と変わらないためここでは省略します。Clspv の SPIR-V を実行する上で気をつける点がいくつかあり、詳細については OpenCL C 1.2 Language on Vulkan を参照して下さい。ここでは大事なものをいくつかピックアップして説明します。
DescriptorSet
Clspv でビルドした OpenCL カーネルの引数は、 Vulkan では DescriptorSet にバインドしたバッファーに対応します。基本的に OpenCL カーネルの引数の順番と DescriptorSet にバインドしたバッファーの順番は対応しています。カーネル引数とバッファーの具体的な対応関係を確認したい場合は、Clspv で Descriptor map を出力できます。Descriptor map は Clspv でビルドする際に --descriptorma=<filename>を付けることで出力できます。例えば、 前節の vulkan_clspv_test1.cl の場合は、
% clspv --c++ --inline-entry-points --descriptormap=vulkan_clspv_test1.csv -o=vulkan_clspv_test1.spv vulkan_clspv_test1.cl
とすることで vulkan_clspv_test1.csv というファイル名のCSV形式の Descriptor map が出力されます。CSVファイルの内容は、
kernel,testSummation,arg,value,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer
kernel,testSummation,arg,outputs,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer
のようになります。各行はそれぞれひとつの引数についての情報が載っています。行の内容は、フィールド名と値が交互に記述されています。行の内容は、
kernel: OpenCLカーネル名arg: 引数名argOrdinal: 引数のポジションdescriptorSet: DescriptorSet のインデックスbinding: 使用するバインドされたバッファーのポジションoffset: バッファーにアクセスする際のアドレスのオフセットargKind: バッファの種類
となります。この、binding の部分を見ればカーネルの引数がどのバッファーに対応するかわかります。
Work-Group Size
OpenCL では Work-Group のサイズは clEnqueueNDRangeKernel を呼ぶ時に引数として指定していました。Clspv でビルドしたカーネルの Work-Group サイズを指定する場合は、reqd_work_group_size を使ってカーネルの属性として指定するか、Specialization Constants を設定して指定する方法があります。Vulkan を動かす GPU が固定な場合は reqd_work_group_size を使用してハードコードする方法で良いですが、動かすGPUが複数あり、GPUによって Work-Group サイズを変える場合は Specialization constants で指定する方法が良いです。Specialization Constants は Compute Shader のパイプラインを作成する時に設定します。
std::array<uint, 3> work_group_size{{x, y, z}}; // X、Y、Z次元の Work-Group サイズ
// X、Y、Zの Work-Group サイズを設定するため、
// それぞれの specialization constant の ID やデータサイズを設定する
std::array<vk::SpecializationMapEntry, 3> entries;
for (std::size_t i = 0; i < entries.size(); ++i) {
entries[i].constantID = static_cast<uint>(i);
entries[i].offset = static_cast<uint>(i * sizeof(uint));
entries[i].size = sizeof(uint);
}
const vk::SpecializationInfo info{3, entries.data(), 3 * sizeof(uint), work_group_size.data()};
// Work-Group サイズの情報を含めて パイプラインを作成する
vk::ShaderModule module = /* Clspv でビルドした SPIR-V を読み込んだモジュール */
const char* kernel_name = /* OpenCL カーネル名 */
const vk::PipelineShaderStageCreateInfo shader_stage_create_info{
vk::PipelineShaderStageCreateFlags{},
vk::ShaderStageFlagBits::eCompute,
module,
kernel_name,
&info};
vk::PipelineLayout pipeline_layout = /* 作成した PipelineLayout */
const vk::ComputePipelineCreateInfo create_info{
vk::PipelineCreateFlags{},
shader_stage_create_info,
pipeline_layout};
vk::Device device = /* 作成した Device */
vk::Pipeline pipeline = device.createComputePipelines(vk::PipelineCache{}, create_info);
例えば、Work-Group の合計サイズを 64 で設定する場合、1次元の Work-Group にする場合は、X=64、Y=1、Z=1 (X * Y * Z = 64) と指定すると、OpenCL 上では、
get_local_size(0); // X = 64
get_local_size(1); // Y = 1
get_local_size(2); // Z = 1
と反映されます。2次元の Work-Group にする場合は、X=8、Y=8、Z=1 (X * Y * Z = 64) と指定すると、OpenCL 上では、
get_local_size(0); // X = 8
get_local_size(1); // Y = 8
get_local_size(2); // Z = 1
と反映されます。
おわりに
この記事では Vulkan SDK と Clspv の導入を簡単に説明しました。Clspv の登場によって OpenCL コードを Vulkan で動かすことが可能になりました。OpenCL は C言語に慣れている方であれば書きやすい言語だと思いますが、OpenCL ドライバのバージョンが 1.2 で止まっていたりと実行環境周りに不安を覚えている方もいると思います。そのような場合は、 Clspv を使って Vulkan を始めてみるのも面白いと思います。
Clspv はまだ開発中のプロトタイプであるため色々と不具合が見つかるかもしれません。その時は github で issue を作って報告してあげて下さい。
