このブログは、株式会社フィックスターズのエンジニアが、あらゆるテーマについて自由に書いているブログです。
OpenCLを書いている時、ホストとデバイスのコードが完全に分離していて連携しづらくてツライ・・・といった経験ありませんか?AMDのGPUに限るならHIPといった選択肢もあるのですが、OpenCLであってほしい場面はたくさんあります。
そんなところで、実は、OpenCLを単一ソースで書けるKhronosの規格にSYCL(シクル、と発音します)というものがあります。
この記事ではこの規格の実装をためしにつかってみました。付属してるサンプルなどを動作させるところまでの解説と、書いてみたコードの紹介をします。
Khronosグループによると、単一ソースのモダンなC++(具体的にはC++11)でOpenCLデバイスを扱うための規格です。バージョン1.2と2.2の暫定版が策定されています。詳しくはKhronosやCodePlayのサイトに書いてありますので詳しく知りたい方はお読みください。
KhronosによるとSYCLの実装は、CodePlayによるComputeCPPとGitHubで開発されているtriSYCLの二つがあるそうです。今回はこの二つの実装をそれぞれ動かしてみました。
今回使った環境は特に明言しない限り以下の通りです
ほかの環境で試すときは適宜読み替えてください。OpenCLが使える新しめなx86_64のLinux環境であればなんとかなるかもしれません。
ComputeCPPはCodePlayが公開しているクローズドソースなSYCL1.2の実装(注: Khronosのサイトでは2.2を実装してあるかのように書いてありますが、CodePlayのサイトを見渡す限り1.2です)で、Community Editionのベータ版が使えます。ライセンスは実行ファイル等と一緒に圧縮されているdoc/LICENSE.TXT
に書いてあります。
この実装の動作としてはComputeCPPはSYCLのコードを一旦LLVMIRに変換し、それをSPIR1.2を使ってGPUで動かしています。したがって、OpenCLデバイスがKhronos拡張のcl_khr_spirに対応している必要があります。
まず、公式サイトでアカウントを作成し、使ってるOSにあったバージョンのComputeCPPをダウンロードし、解凍します。この記事ではこれ以降、解凍したフォルダが/opt/ComputeCPP-CE
にあることにします。
なお、ComputeCPPはCentOSとUbuntu 14.04、Ubuntu16.04のそれぞれの64bit版に対応しているバイナリを提供しています。
CodePlayがComputeCPPのためのユーティリティとドキュメントをgithubで公開しています。この中にサンプルコードもあります。また、このリポジトリはApache Licenseでライセンスされています。
$ git clone https://github.com/codeplaysoftware/computecpp-sdk
サンプルは以下の手順でビルドできます。(筆者の環境ではfind_package(OpenCL REQUIRED)
がうまく動かなかったため-DOpenCL_LIBRARY
をつけています)
$ cd computecpp-sdk/samples/
$ mkdir build && cd build
$ cmake -DCMAKE_BUILD_TYPE=Release -DCOMPUTECPP_PACKAGE_ROOT_DIR=/opt/ComputeCpp-CE/ -DOpenCL_LIBRARY=/opt/AMDAPPSDK-3.0/lib/x86_64/libOpenCL.so ../
$ make
しかし、この方法で作成したバイナリは筆者の環境だと実行時にSegmentation Faultになってしまいます。(どうやらNVIDIAのプロプライエタリドライバを使った環境ではSegmantation Faultにならないこともあるようです)
しかし、以下のように直接ビルドするとSegmentation Faultにならないようです。ここではsamples/matrix_multiply.cpp
をビルドするものとします。
cd computecpp-sdk/samples/matrix_multiply
/opt/ComputeCpp-CE/bin/compute++ -std=c++11 -O3 -I/opt/ComputeCpp-CE/include -c -sycl \
-emit-llvm -I/opt/AMDAPPSDK-3.0/include/ -no-serial-memop matrix_multiply.cpp
/opt/ComputeCpp-CE/bin/compute++ -std=c++11 -O3 -I/opt/AMDAPPSDK-3.0/include/ \
-I/opt/ComputeCpp-CE/include/ -L/opt/AMDAPPSDK-3.0/lib/x86_64 -lOpenCL -lpthread \
-fopenmp /opt/ComputeCpp-CE/lib/libComputeCpp.so -include matrix_multiply.sycl matrix_multiply.cpp
-Iや-Lは環境や設定に応じて適宜書き換えてください。二行目でmatrix_multiply.sycl
というファイルを生成します。この.syclにSPIRが含まれているようです。
実行すると以下のように出力が得られました。
$ ./a.out 8192 sycl
Input matrix
***** SYCL
The Device Max Work Group Size is : 256
The order is : 8192
The blockSize is : 8
SYCL: Time: 4729
GFLOPs: 232.504
Output
Success
triSYCLはオープンソースで開発されている実装で、ComputeCPPと違いインクルードするだけで使うことができます。ライセンスはLLVM/Clangと同じようです。
OpenCLデバイスでコードを走らせることはできず、普通に書くとOpenMPでSYCLのカーネルを並列実行することができるだけのようです。boost::computeのkernelを使ったりするinterperability modeというものを使うことでOpenCLデバイスを使うこともできますが、SYCLで書くことはできないようなのでここでは割愛します。
GitHub からcloneするだけです。
git clone https://github.com/triSYCL/triSYCL.git
triSYCLがboost::compute依存しているので、インストールしておきます。
sudo apt install libcompute-dev
triSYCLに付属してるテストをビルドします。boost::computeが/usr/include/computeにあったので、指定を追加しています。
cd triSYCL/tests/
echo "CXXFLAGS+=-I/usr/include/compute" >> Makefile
make -j8
ビルドには数分程度かかるかもしれません。また、ビルド中に大量のwarningが出ます。
make run
とするとテストが実行されます。最後の行に
**** no errors detected
と出ていれば成功です。
先ほど動かしたComputeCPPのサンプルとtriSYCLのテストはそれぞれのライブラリでしかビルドできないようなので、どちらでもビルド & 実行できるプログラムを書いてみました。
#include<iostream>
#include<CL/sycl.hpp>
constexpr std::size_t NDIM=100;
// 初期化
template<typename T = std::int64_t>
void set_value(T vec1[],T vec2[],std::size_t size)
{
for (auto i = decltype(NDIM)(0); i < size; i++)
{
vec1[i] = -i;
vec2[i] = 2*i;
}
}
// 計算本体
template<typename T = std::int64_t>
void calc(T vec1[],T vec2[],std::size_t size)
{
auto dim = size;
// myQueueというqueueを作る。
// host_selectorを指定しておくとhostで実行される
// cpu_selectorやgpu_selecotr等もあるがtriSYCLでは動かない。
cl::sycl::queue myQueue(cl::sycl::host_selector {});
// T型の長さdimの配列であるvec1とvec2をVec1,Vec2というbufferごしに扱えるようにする。
cl::sycl::buffer<T> Vec1 {vec1,dim};
cl::sycl::buffer<T> Vec2 {vec2,dim};
// myQueueにラムダ式をsumit。
// このラムダ式内でデータへのアクセスの制御とOpenCLで実行することを書く。
// cl::sycl::handlerはkernelで使っているdeviceにあるhost dataに関する、
// 正確な制御のために必要とされるすべての仕事をこなすためにある。queueごとにhandlerがある。
// cghとはcommand group handlerのこと
myQueue.submit([&Vec1,&Vec2,&dim](cl::sycl::handler &cgh)
{
// bufferにaccessするためにaccessorを作る
// v1は読み書き可能
// v2は読み込みのみ可能
auto v1 = Vec1.template get_access<cl::sycl::access::mode::read_write>(cgh);
auto v2 = Vec2.template get_access<cl::sycl::access::mode::read>(cgh);
// この行以降でOpenCLデバイスで実行すること(kernel)を書く。
// cgh.parallel_forを使うと並列実行される。
// cl::sycl::range<1>(dim)で1次元でdim要素であることを示し、
// cl::sycl::id<1> indexで1次元で0以上dim未満の範囲の数が
// ラムダ式の引数としてその後のコードに渡され、それぞれ独立に並列実行される。
cgh.parallel_for<class vectoradd>(cl::sycl::range<1>(dim),[v1,v2,&dim](cl::sycl::id<1> index)
{
// bufferへは普通の配列のように[]を使える
// bufferに書き込めば元の配列にも反映される
v1[index] = v1[index] + v2[index];
});
});
}
// 出力
template<typename T = std::int64_t>
void output_value(T vec1[],std::size_t size)
{
for (auto i = decltype(size)(0); i < size; i++)
{
std::cout << vec1[i] << std::endl;
}
}
int main(void)
{
int64_t vec1[NDIM],vec2[NDIM];
set_value<>(vec1,vec2,NDIM);
calc<>(vec1,vec2,NDIM);
output_value<>(vec1,NDIM);
return 0;
}
triSYCLでは
g++ vectoradd.cpp -o vectoradd_sycl -fopenmp -g -isystem/opt/AMDAPPSDK-3.0/include \
-isystem/opt/triSYCL/include -L/opt/AMDAPPSDK-3.0/lib/x86_64 -lOpenCL -lpthread -std=c++1y -O3
とするとビルドでき(現状ではOpenCLをリンクせずとも、また-std=c++14でも動きますが、ここで提示したオプションは念のためです)、
./vectoradd_sycl
とすると実行できます。
ComputeCPPでは
/opt/ComputeCpp-CE/bin/compute++ -std=c++11 \
-I/home/yuki_ito/ComputeCpp-CE-0.2.0-Linux/include \
-I/opt/AMDAPPSDK-3.0/include/ -c -sycl -emit-llvm vectoradd.cpp
g++ -I/opt/ComputeCpp-CE/include -L/opt/ComputeCpp-CE/lib -lComputeCpp \
-g -I/opt/AMDAPPSDK-3.0/include/ -L/opt/AMDAPPSDK-3.0/lib/x86_64 -lOpenCL \
-include vectoradd.sycl vectoradd.cpp -o vectoradd_computecpp
とするとビルドでき、実行のためにはlibComputeCpp.soが必要なので、
export LD_PRELOAD=~/ComputeCpp-CE-0.2.0-Linux/lib/libComputeCpp.so
./vectoradd_computecpp
としてやれば動きます。
二つの実装について一言でまとめると、
といったところでしょうか。調べてる過程でtriSYCLもComputeCPPのどちらも未実装な部分が見られたり、動作が怪しい部分が見つかったりしたものの、活発に開発されているようなので筆者としてはこれからに期待しています。
今回の記事は動かすところまでだったので、性能評価などはおこなっていませんが、機会があれば性能評価を行った記事も書きたいですね。
ルネサス社が、ComputeCPPの開発元であるCodeplayと共同で、SYCLを使ってADAS(先進運転システム)に取り組むそうです。SYCLの今後が楽しみですね!
Renesas Electronics and Codeplay Collaborate on #OpenCL and #SYCL for ADAS Solutions https://t.co/hNwyMAmaIA
— OpenCL API (@openclapi) September 12, 2017
Renesas+Codeplayの中の人です。がんばるので、ぜひ応援してくださいね。