Article

2019年2月26日
みなさん、今日は。今日も元気にCUDAやってますか?

ソフトウェア高速化の技術者ならCUDAぐらいできて当然になって久しい世の中(フィックスターズ社内)(※個人の感想です)ですが、「より高品質で効率的な開発を」という要求は今も変わらず続いています。というか、CUDAが当たり前になって実業務アプリに近いところでも応用が広がっている昨今だからこそ、そのような需要が高まっていると感じます。

そんな中で、今回は1つ小ネタを紹介します(n番煎じかもしれませんが)。それはタイトルの通り、CUDAのデバイスメモリを、C++のスマートポインタに扱わせる方法です。

TL;DR

どんな感じかは、コードを見てもらったほうが早いかと思います。

// このコードは、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};
	}
}

CC0

解説

CUDAのデバイスメモリを利用するには、C APIであるcudaMalloc()し、使い終わったらcudaFree()します。
生のC言語を使うならそれでいいのですが、大規模で複雑になるとC言語では開発コストが高すぎます。フィックスターズの業務でも、組み込み機器やSoCなどの貧弱な環境でなく、GPUが使えるような環境ではお客様から提供されるコードもC++であることが多く、C++を適切に用いてCUDAを扱うことが求められます。

そして、現代の(つまり、11より前の石器時代みたいなのではない)C++では、このようなメモリ・リソースは生で扱わずスマートポインタに管理を任せることが常識です。ですので、CUDAのデバイスメモリも当然スマートポインタ(std::unique_ptrstd::shared_ptr)で扱う必要が出てきます。

これを実現するのが、今回紹介する方法です。使い方は簡単。通常std::unique_ptrstd::make_uniqueを使うところで、代わりにcuda::unique_ptrcuda::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
{
	void check_error(const ::cudaError_t e, decltype(__FILE__) f, decltype(__LINE__) 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__))

CC0

まとめ

CUDAのデバイスメモリ管理もスマートポインタにまかせて、楽しいCUDAプログラミングを!

※今回紹介したコードは、wandbox上でclang++とg++のC++11で動作確認していますが、CC0のライセンス条文にもある通り動作結果については無保証であり、著者およびフィックスターズは、これを利用したことによるいかなる損害に関する責任も負いません。

Tags

About Author

YOSHIFUJI Naoki

yoshifujiです。計算力学的なプログラムを高速化することが得意です。プログラミング自体はチョットダケワカリマス。 Twitter: https://twitter.com/LWisteria

Leave a Comment

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

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

Recent Comments

Social Media