SYCLを使ってOpenCLを単一ソースで書いてみる

2017年8月17日

OpenCLを書いている時、ホストとデバイスのコードが完全に分離していて連携しづらくてツライ・・・といった経験ありませんか?AMDのGPUに限るならHIPといった選択肢もあるのですが、OpenCLであってほしい場面はたくさんあります。

そんなところで、実は、OpenCLを単一ソースで書けるKhronosの規格にSYCL(シクル、と発音します)というものがあります。

この記事ではこの規格の実装をためしにつかってみました。付属してるサンプルなどを動作させるところまでの解説と、書いてみたコードの紹介をします。

SYCLについて

Khronosグループによると、単一ソースのモダンなC++(具体的にはC++11)でOpenCLデバイスを扱うための規格です。バージョン1.2と2.2の暫定版が策定されています。詳しくはKhronosCodePlayのサイトに書いてありますので詳しく知りたい方はお読みください。

KhronosによるとSYCLの実装は、CodePlayによるComputeCPPとGitHubで開発されているtriSYCLの二つがあるそうです。今回はこの二つの実装をそれぞれ動かしてみました。

実験環境

今回使った環境は特に明言しない限り以下の通りです

  • OS : Ubuntu 16.04
  • GPU : AMD Radeon R9 Fury X
  • ドライバ : AMDGPU-Pro Driver Version 17.10 for Ubuntu 16.04.2
  • SDK : AMD APP SDK v3.0

ほかの環境で試すときは適宜読み替えてください。OpenCLが使える新しめなx86_64のLinux環境であればなんとかなるかもしれません。

ComputeCPP

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に対応している必要があります。

準備

ComputeCPPの入手

まず、公式サイトでアカウントを作成し、使ってるOSにあったバージョンのComputeCPPをダウンロードし、解凍します。この記事ではこれ以降、解凍したフォルダが/opt/ComputeCPP-CEにあることにします。

なお、ComputeCPPはCentOSとUbuntu 14.04、Ubuntu16.04のそれぞれの64bit版に対応しているバイナリを提供しています。

ComputeCpp-SDKの入手

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

triSYCLはオープンソースで開発されている実装で、ComputeCPPと違いインクルードするだけで使うことができます。ライセンスはLLVM/Clangと同じようです。

OpenCLデバイスでコードを走らせることはできず、普通に書くとOpenMPでSYCLのカーネルを並列実行することができるだけのようです。boost::computeのkernelを使ったりするinterperability modeというものを使うことでOpenCLデバイスを使うこともできますが、SYCLで書くことはできないようなのでここでは割愛します。

準備

triSYCLの入手

GitHub からcloneするだけです。

git clone https://github.com/triSYCL/triSYCL.git

boost::compute

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

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

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
OpenMPの使える環境であれば手軽にSYCLを書くことができるが、並列処理を書きたいだけであれば直接OpenMPで書いたほうが良いかと思います。
ComputeCPP
サイトに登録したり、専用のビルド手順を踏む必要があるものの、SYCLで書いたものをOpenCLで動かすことができる。

といったところでしょうか。調べてる過程でtriSYCLもComputeCPPのどちらも未実装な部分が見られたり、動作が怪しい部分が見つかったりしたものの、活発に開発されているようなので筆者としてはこれからに期待しています。

今回の記事は動かすところまでだったので、性能評価などはおこなっていませんが、機会があれば性能評価を行った記事も書きたいですね。

2017/09/29追記

ルネサス社が、ComputeCPPの開発元であるCodeplayと共同で、SYCLを使ってADAS(先進運転システム)に取り組むそうです。SYCLの今後が楽しみですね!

Tags

About Author

yuki.ito

1件のコメント

  • Renesas+Codeplayの中の人です。がんばるので、ぜひ応援してくださいね。

Leave a Comment

メールアドレスが公開されることはありません。 が付いている欄は必須項目です

このサイトはスパムを低減するために Akismet を使っています。コメントデータの処理方法の詳細はこちらをご覧ください

Recent Comments

Social Media