[C++] 標準化された並列アルゴリズムを用いたCUDA GPUの利用
2回ほどOpenACCを紹介しました*が、今回はより簡単に並列化できる手法としてC++17規格で標準化されたC++ Parallel Algorithmsを紹介します。
NVIDIA HPC SDKで利用できるC++コンパイラのnvc++では、CPUとGPUで並列化された各アルゴリズムを提供しており、すでにC++で計算コードを書いている場合、OpenACCよりも簡単にGPU化の恩恵を得られそうです。
この記事では並列アルゴリズムの概要を簡単に紹介し、サンプルコードとコンパイル方法を記載します。また、並列アルゴリズムの代替として利用可能なCUDA Thrust Libraryも併せて紹介します。
実行環境は以下を想定しています。
- CUDA SDK 12.1
- CUDA Driver 531.18 (CUDA 12.1)
- NVIDIA HPC SDK 23.1
- GCC 11.3.0
C++並列アルゴリズム
C++では<algorithm>でSTLコンテナに対する様々なアルゴリズムを提供していますが、C++17では<algorithm>の大部分と<numeric>に追加されたreduceおよびtransform_reduceに並列化されたバージョンが実装されています。
各アルゴリズムの第一引数に、Execution policyを受け取る実装が追加され、プログラマはポリシーを与えることで言語実装 (コンパイラ) に並列化を許可します。
言語仕様上は並列化が可能であることを示すものなので、実際に並列化するかは言語実装が決定します。nvc++コンパイラでは、マルチコアCPUまたはCUDA GPUでの並列実行を試みます。
std::transformを用いたfloat型のスカラとベクトルの積SAXPYの実装例が以下の通りです。ラムダ式も利用できます。
#include <iostream>
#include <vector>
#include <execution>
#include <algorithm>
int main()
{
constexpr std::size_t n = 12;
std::vector x(n, 1.0f), y(n, 0.5f);
float alpha = 1.25f;
// Parallel SAXPY
std::transform(std::execution::par_unseq,
std::begin(x), std::end(x), std::begin(y), std::begin(y),
[=] (float x, float y) {
return alpha * x + y;
});
// Sequential execution
std::for_each(std::cbegin(y), std::cend(y),
[] (float v) {
std::cout << v << std::endl;
});
}
std::execution::par_unseqポリシーはnvc++でサポートするすべてのCUDA GPUでの実行を許可します。std::execution::parポリシーもCUDA GPUでの実行を許可しますが、Tesla P100などPascal世代のGPUでは実行されないようです。
上記のサンプルコードで実際に並列化されたバージョンを使うためには、-stdparコンパイルオプションを指定するだけです。デフォルトではGPUで並列化されますが、-stdpar=multicoreを指定するとマルチコアCPUで並列化されます。
# Parallel algorithms on CUDA GPU
$ nvc++ -std=c++17 -stdpar sample.cc
$ nvc++ -std=c++17 -stdpar=gpu sample.cc
# Parallel algorithms on Multi-core CPU
$ nvc++ -std=c++17 -stdpar=multicore sample.cc
$ ./a.out
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
NVIDIA HPC SDKのドキュメントを見ると、nvc++での並列アルゴリズムの実装は以前紹介したCUD A unified memoryを活用してCPU-GPU間のデータ転送を行っているようです。したがって、Unified memoryのメリット・デメリットがそのまま適用されます。
制限
C++標準アルゴリズムでは様々なタイプのイテレータをサポートしていますが、GPUを用いた並列アルゴリズムはRandom Access Iteratorのみがサポートされます。簡単にはポインタと同等の機能を持つイテレータの必要がある、ということです。
また例外はサポートされません。例外コードはコンパイルできるようですが、catchは無視されるのと、throwが起きるとGPU処理が停止するようです。
複雑なループを並列化するには
<algorithm>はイテレータを渡して処理を行うため、通常は1個または2個のデータ構造への処理が想定されています。
定型化された処理を行うには向いていますが、3つ以上のデータ構造を必要とするような場合には少し工夫が必要です。
NVIDIA HPC SDKのドキュメントにはより実践的なサンプルとして、LULESHというミニアプリをParallel Algorithm実装について述べられています。
LULESHの実装ではcounting_iteratorというイテレータを実装し、std::for_each_nと組み合わせて並列化したfor文を実現しています。
constexpr int n = 12;
std::vector x(n, 1.0f), y(n, 0.5f);
float alpha = 1.25f;
float const* px = x.data();
float * py = y.data();
std::for_each_n(std::execution::par_unseq, counting_iterator(0), n,
[=](int i) {
py[i] = alpha * px[i] + py[i];
});
より複雑なループであればOpenACCで書くのも選択肢とありますが、上記はあくまでもC++言語仕様にしたがって実装されていますのでnvc++以外の仕様準拠したコンパイラで同様に動作することが保証されます。(先に説明したとおり並列化されるかは実装依存ですが)
CUDA Thrust Library
C++の<algorithm>-likeな並列処理ライブラリとして、CUDA Thrust libraryがあります。
std::vector<T> class template相当のコンテナと、C++ Algorithmに近い並列処理APIを提供します。
並列化は同じくCPUおよびCUDA GPUをサポートしており、コンテナ型やPolicyで実行環境を制御できます。
また、CUDA SDKやNVIDIA HPC SDKにはThrustが同梱されているため、CUDAがすでに導入されている環境では追加のインストールが不要です。
Thrustで同等の処理となるように先ほどのサンプルコードを実装すると、以下のようになります。
#include <iostream>
#include <thrust/universal_vector.h>
#include <thrust/execution_policy.h>
#include <thrust/transform.h>
int main()
{
constexpr std::size_t n = 12;
thrust::universal_vector x(n, 1.0f), y(n, 0.5f); // Unified memory
float alpha = 1.25;
// Run on GPU
thrust::transform(thrust::device, x.begin(), x.end(), y.begin(), y.begin(),
[=] __device__ (float x, float y) {
return alpha * x + y;
});
// Run on CPU
thrust::for_each(thrust::host, y.begin(), y.end(),
[] (float v) {
std::cout << v << std::endl;
});
}
並列アルゴリズムでは動的確保したメモリは自動的にUnified memoryで確保されますが、Thrustの場合はそれぞれCPUメモリとGPUメモリ用、Unified memory用の専用のvector型を用います。今回はUnified memory用のthrust::universal_vectorを利用します。
また、ThrustでもExecution policy相当の機能があり、thrust::device, thrust::hostでそれぞれGPUとCPUどちらで実行するかを明示できます。
こちらはCUDAコンパイラのnvccに、CUDAカーネルとしてラムダ式を利用するために--extended-lambdaオプションをつければ実行できます。
$ nvcc -std=c++17 --extended-lambda main.cu
$ ./a.out
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
1.75
まとめ
簡単に標準化された並列アルゴリズムの使い方を紹介しました。言語標準機能なので学習コストはOpenACCよりも低いと考えられ、プログラムによってはOpenACCよりも低コストでGPUを用いた並列化が実現可能です。また言語拡張であるOpenACCに対しC++ Parallel Algorithmは言語仕様のため、可搬性が重要な場合は有効な選択肢といえます。
一方でメモリは自動的にUnified memoryとして管理されるようなので、明示的なデータ管理を行いたい場合にはThrustライブラリが代替として有効です。
同機能はC++の言語仕様ですが、C++ Algorithmは計算範囲を [begin, end) で指定できれば使えるため、元がC言語で書かれたコードでも十分に使えそうです。こちらも検討の余地があるかと思います。
HPCWORLD補足
- 過去のOpenAC関連コラムは以下で参照できます。
- 近々に本コラムの続編を掲載予定です。