このブログは、株式会社フィックスターズのエンジニアが、あらゆるテーマについて自由に書いているブログです。
ソフトウェア高速化の技術者ならCUDAぐらいできて当然になって久しい世の中(フィックスターズ社内)(※個人の感想です)ですが、「より高品質で効率的な開発を」という要求は今も変わらず続いています。というか、CUDAが当たり前になって実業務アプリに近いところでも応用が広がっている昨今だからこそ、そのような需要が高まっていると感じます。
そんな中で、今回は1つ小ネタを紹介します(n番煎じかもしれませんが)。それはタイトルの通り、CUDAのデバイスメモリを、C++のスマートポインタに扱わせる方法です。
どんな感じかは、コードを見てもらったほうが早いかと思います。
// このコードは、CC0 1.0 全世界(パブリックドメイン)としますので、ご自由にコピペしてお使いください https://creativecommons.org/publicdomain/zero/1.0/deed.ja
// This code is licensed under CC0 1.0 Universal (Public Domain). You can use this without any limitation. https://creativecommons.org/publicdomain/zero/1.0/deed.en
#include <memory>
#include <type_traits>
#include <cuda_runtime_api.h>
namespace cuda
{
struct deleter
{
void operator()(void* p) const
{
CHECK_CUDA_ERROR(::cudaFree(p));
}
};
template<typename T>
using unique_ptr = std::unique_ptr<T, deleter>;
// auto array = cuda::make_unique<float[]>(n);
// ::cudaMemcpy(array.get(), src_array, sizeof(float)*n, ::cudaMemcpyHostToDevice);
template<typename T>
typename std::enable_if<std::is_array<T>::value, cuda::unique_ptr<T>>::type make_unique(const std::size_t n)
{
using U = typename std::remove_extent<T>::type;
U* p;
CHECK_CUDA_ERROR(::cudaMalloc(reinterpret_cast<void**>(&p), sizeof(U) * n));
return cuda::unique_ptr<T>{p};
}
// auto value = cuda::make_unique<my_class>();
// ::cudaMemcpy(value.get(), src_value, sizeof(my_class), ::cudaMemcpyHostToDevice);
template<typename T>
cuda::unique_ptr<T> make_unique()
{
T* p;
CHECK_CUDA_ERROR(::cudaMalloc(reinterpret_cast<void**>(&p), sizeof(T)));
return cuda::unique_ptr<T>{p};
}
}
CUDAのデバイスメモリを利用するには、C APIであるcudaMalloc()
し、使い終わったらcudaFree()
します。
生のC言語を使うならそれでいいのですが、大規模で複雑になるとC言語では開発コストが高すぎます。フィックスターズの業務でも、組み込み機器やSoCなどの貧弱な環境でなく、GPUが使えるような環境ではお客様から提供されるコードもC++であることが多く、C++を適切に用いてCUDAを扱うことが求められます。
そして、現代の(つまり、11より前の石器時代みたいなのではない)C++では、このようなメモリ・リソースは生で扱わずスマートポインタに管理を任せることが常識です。ですので、CUDAのデバイスメモリも当然スマートポインタ(std::unique_ptr
やstd::shared_ptr
)で扱う必要が出てきます。
これを実現するのが、今回紹介する方法です。使い方は簡単。通常std::unique_ptr
やstd::make_unique
を使うところで、代わりにcuda::unique_ptr
やcuda::make_unique
を使うだけです(名前空間を変えれば良い)。
make_uniqueについては、標準ライブラリに合わせて配列版と非配列版が用意されています。通常は配列版を使うことが多いだろうと思います。
ただし注意すべきは、非配列版は、標準ライブラリにあるコンストラクタの引数を取るインターフェースは用意していないことです。これは、CUDAではデバイス上に直接コンストラクトできないのでホストメモリ側で一度インスタンスを生成する必要がありますが、非配列版で扱うオブジェクトの大きさが不明(一般的には大きいと予想される)ため、効率があえて悪化すると考えられるからです。
配列版や通常のCUDAプログラミングと同様に、確保したポインタへcudaMemcpy()
などで転送してください。
間違って配列版を(1つの整数引数をとる)コンストラクタ版と混同しないように、引数を取るmake_uniqueはSFINAEによって配列型でないと呼べないようにしてあります。
また、ここではunique_ptrしか紹介しませんでしたが、shared_ptrもほぼ同様なので割愛します。
CHECK_CUDA_ERROR()
上記コードに含まれているCHECK_CUDA_ERROR()
はC APIの戻り値(cudaError_t
)を確認するもので、ご自身で利用しやすいものを用意していただければ良いと思います。
参考までに、私が普段使っているものを載せておきます。
// このコードは、CC0 1.0 全世界(パブリックドメイン)としますので、ご自由にコピペしてお使いください https://creativecommons.org/publicdomain/zero/1.0/deed.ja
// This code is licensed under CC0 1.0 Universal (Public Domain). You can use this without any limitation. https://creativecommons.org/publicdomain/zero/1.0/deed.en
#include <stdexcept>
#include <sstream>
#include <cuda_runtime_api.h>
namespace cuda
{
template<typename F, typename N>
void check_error(const ::cudaError_t e, F&& f, N&& n)
{
if(e != ::cudaSuccess)
{
std::stringstream s;
s << ::cudaGetErrorName(e) << " (" << e << ")@" << f << "#L" << n << ": " << ::cudaGetErrorString(e);
throw std::runtime_error{s.str()};
}
}
}
#define CHECK_CUDA_ERROR(e) (cuda::check_error(e, __FILE__, __LINE__))
CUDAのデバイスメモリ管理もスマートポインタにまかせて、楽しいCUDAプログラミングを!
※今回紹介したコードは、wandbox上でclang++とg++のC++11で動作確認していますが、CC0のライセンス条文にもある通り動作結果については無保証であり、著者およびフィックスターズは、これを利用したことによるいかなる損害に関する責任も負いません。
`check_error` の引数 `f` の型、 `std::decay_t` にしないと `check_error` がおいてあるファイルと同じ長さのファイル名を持つファイル内でしか呼べなくないですか