[C++] OpenACCとCUDA (Thrust library) を連携する
業務や微力ながらメンターとしてお手伝いしているGPUミニキャンプなどで、OpenACCとCUDAを組み合わせる方法について尋ねられることが何度かあり、今回1例ですがご紹介します。
何度かOpenACCのメリットについて紹介してきました。OpenACCは物理シミュレーションで頻出するような計算量が大きく並列化しやすいループを、GPUで簡単に並列化するのに適しています。
一方で並列化するときにスレッド間でデータのやり取りが必要であったり、コンパイラがうまく並列化を行えなかったりすると途端にOpenACCの利用が難しくなることがあります。特に、細かい最適化が必要になるケースでかゆいところに手が届きにくいようです。 (私もCUDAならこの機能が使えるのに、ということはありました)
より高度な最適化を行いたい場合、CUDAを直接使った実装は可能です。しかし、メンテナンス性は悪くなりますのでプログラム全体をCUDA化することは現実的ではありません。
またBLASやFFTいった頻出する計算はNVIDIAから最適化されたCUDAライブラリが提供されていて、OpenACCを使って自前で書くよりも性能も生産性も高いため、OpenACCコードと組み合わせて活用できると嬉しいです。
このような現実的な問題に対処するため、OpenACC Specificationにはポインタの相互変換や、CUDAのContextやStreamを取得するためのAPIが提供されています。
この機能により例えばOpenACCで実装されたコードの一部だけをCUDA化したり、CUDAライブラリをOpenACCプログラムに組み込むといったことが可能になります。
この記事ではC++ template機能を使ってメモリ管理や並列処理を提供するThrust libraryとの連携を例に、OpenACCからCUDAを、CUDAからOpenACCを、相互に呼び出す方法について紹介します。
Thrustについては前回記事を参照いただければと思いますが、簡単にはC++ Parallel Algorithmに近いインターフェイスを提供するCUDAライブラリで、配列全体に対する定型化された処理についてはOpenACCで書くよりもThrustで実装したほうが簡潔かつ高速であると期待されます。
実行環境は以下を想定しています。
- CUDA SDK 12.1
- CUDA Driver 531.18 (CUDA 12.1)
- NVIDIA HPC SDK 23.1
- GCC 11.3.0
OpenACCとThrustの連携
OpenACCとThrustの連携を具体的に考えると、以下のようなメリットがあります。
- もともとOpenACCで書いていて、OpenACCのメモリをThrustで処理する
- もともとThrustで書いていて、ThrustのメモリをOpenACCで処理する
- 複数のデータ構造を組み合わせた計算はThrustでは難しいため、OpenACCで実装したほうが生産性が高い
- CUDAの実装はメンテナンスが困難になるため、CUDAライブラリを使っていてOpenACCと連携することも十分にあり得る
プログラムのビルドは、nvc++ (C++コンパイラ) とnvcc (CUDA C++コンパイラ) で各ソースファイルをビルドし、nvc++でリンクして実行ファイルを生成します。
以下のような流れとなり、2つのコンパイラを組み合わせる必要があるため少し大変ですが、nvc++とnvccの場合はリンクが必要なOpenACCとCUDAライブラリをオプション1つで済ませることができるので比較的簡単になっています。
# OpenACCコードをコンパイル
nvc++ -acc -c openacc_code.cc
# CUDAコードをコンパイル
nvcc -c cuda_code.cu
# 必要なライブラリをリンクして実行ファイルを生成
# -cuda optionでCUDAのランタイムライブラリをリンクします
nvc++ -acc -cuda -o xtest openacc_code.o cuda_code.o
OpenACCからThrustを呼び出す
先にOpenACCからThrust (CUDA) を呼び出す場合の実装例を紹介します。
OpenACCからThrust (CUDA) を呼び出す場合に要となるのが、acc host_data use_device指示句です。
#pragma acc host_data use_device(x)
{
cuda_native_func(x);
}
use_device(x)は次のOpenACCコードブロック内で、OpenACCが管理するメモリxをdevice native (ここではCUDA) なraw pointerとします。
raw pointerなので、これをそのままCUDAカーネルやAPIに渡すことが可能です。
前回と同様、簡単な例としてSAXPY相当のThrustで実装し、OpenACCから呼び出すプログラムを実装しました。
以下のソースコード内に、必要なコメントを記載しています。
main.cc
#include <iostream>
#include <algorithm>
#include <vector>
// CUDAカーネル
void saxpy(std::size_t n, float a, float const* x, float * y);
int main()
{
constexpr std::size_t n = 12;
std::vector<float> x(n, 1.0f), y(n, 0.0f);
float alpha = 0.5f;
auto x_ptr = x.data();
auto y_ptr = y.data();
/*
std::vector<T> で確保したメモリを
raw pointerにしてOpenACCに渡す
*/
#pragma acc data copyin(x_ptr[0:n]) copy(y_ptr[0:n])
{
/*
host_data use_device でOpenACC pointerを
CUDA native device pointerにする
*/
#pragma acc host_data use_device(x_ptr, y_ptr)
{
saxpy(n, alpha, x_ptr, y_ptr);
}
}
std::for_each(std::cbegin(y), std::cend(y), [](float v) {
std::cout << v << std::endl;
});
}
kernel.cu
#include <thrust/transform.h>
#include <thrust/copy.h>
#include <thrust/execution_policy.h>
// x and y as CUDA device pointers
void saxpy(std::size_t n, float a, float const* x, float * y)
{
/*
この関数ではx, yがどこに確保されたメモリかコンパイラでは判断できない
thrust::device execution-policyをつけることでCUDA kernelとして実行する
*/
thrust::transform(thrust::device, x, x + n, y, y,
/*
CUDAコンパイラの制約で __device__ が必要
またコンパイル時には -extended-lambda オプションが必要
*/
[=] __device__ (float x, float y) {
return a * x + y;
});
}
コンパイルと実行
$ nvc++ -acc -c main.cc
$ nvcc -extended-lambda -c kernel.cu
$ nvc++ -acc -cuda -o xtest main.o kernel.o
$ ./xtest
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
ThrustからOpenACCを呼び出す
次にThrustからOpenACCを呼び出します。NVIDA HPC SDKのOpenACCではメモリをすべてCUDA unified memoryで確保し、CPU-GPU間のデータ転送を自動化することが可能です。
一方すべてUnified memoryになってしまうため、GPU上にだけ置きたいデータがある場合に無駄なコストがかかってしまう、一部のデータ転送だけの制御はできない、といったデメリットもあります。
とはいえCUDA APIを直接利用するとメモリ管理が大変になってしまう、そこでThrustのvector<T>型を使って管理を容易にする、という使い方が考えられます。
ThrustにもUnified memoryが利用できるuniversal_vector<T>型が提供されているので、OpenACC単体ではできなかった「基本Unified memoryで転送を自動化しつつ一部を厳密に制御する」が容易に実現できます。
Thrust (CUDA) からOpenACCを呼び出す際には、acc deviceptr指示句を使用します。
#pragma acc kernels deviceptr(x)
for (std::size_t i = 0U ; i < N ; ++i)
x[i] = i;
deviceptr(x)は次のOpenACCコードブロック内で、xをdevice nativeなraw poitnerとみなし、カーネル内でそのまま利用します。
通常OpenACCカーネル内でポインタを使うとメモリ確保やデータ転送等を行おうとしますが、xはすでにGPU上に確保されたメモリを指しているポインタと定義したため、それらの処理が行われません。
こちらもSAXPY相当の処理をOpenACCで記述し、Thrustで管理したメモリを使って計算するプログラムを実装しました。
以下のソースコード内に、必要なコメントを記載しています。
main.cu
#include <iostream>
#include <algorithm>
#include <thrust/universal_vector.h>
// OpenACCカーネル
void saxpy(std::size_t n, float alpha, float const* x, float * y);
int main()
{
constexpr std::size_t n = 12;
/*
thrust::universal_vector はunified memoryを使った
host/device どちらからでもアクセスできる動的配列
*/
thrust::universal_vector<float> x(n, 1.0f), y(n, 0.0f);
float alpha = 0.5f;
/*
thrust::*_vector からraw pointerを取得するには thrust::raw_pointer_cast を使う
*/
saxpy(n, alpha, thrust::raw_pointer_cast(x.data()), thrust::raw_pointer_cast(y.data()));
/*
unified memoryにより自動的にGPU-CPU間でデータが転送される
*/
std::for_each(std::cbegin(y), std::cend(y), [](float v) {
std::cout << v << std::endl;
});
}
kernel.cc
#include <cstddef>
// x and y as CUDA device pointers
void saxpy(std::size_t n, float alpha, float const* x, float * y)
{
/*
deviceptr でraw pointerをdevice pointerとして扱う
データ転送などはされずにそのままOpenACCカーネルで利用される
*/
#pragma acc kernels deviceptr(x, y)
for (std::size_t i = 0U ; i < n ; ++i)
y[i] = alpha * x[i] + y[i];
}
コンパイルと実行
$ nvcc -c main.cu
$ nvc++ -acc -c kernel.cc
$ nvc++ -acc -cuda -o xtest main.o kernel.o
$ ./xtest
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
0.5
まとめ
OpenACCとThrust (CUDA) の連携をするための機能を紹介し、2つのサンプルを紹介しました。
今回の方法でThrustだけでなく、cuBLASやcuFFTなどNVIDIAが提供するCUDAライブラリを呼び出すことも可能です。
またすでにCUDAで実装したコードがあれば、それを呼び出すことも可能と考えられます。
メモリ管理をあえて行いたいという場合には、CUDAからOpenACCを呼び出す手法も有効と思われます。
その場合はThrustのvector型を使うことで、CUDA APIを直接使うよりも開発効率が高くなると期待されます。
