NanoSTL(a portable and small subset of C++ STL) が CUDA で使えるようになってきました

背景

• CUDA(device 側)でもおてがるぱぱっと時短で std::vector とか使いたい
• 特に性能は気にしない(データのセットアップだったりと, パフォーマンスに影響しないところで使う)
• CPU 実行のコードと統一したい(CUDA 用に書き換えるのめんどかったり, CUDA でうまく動かないときに CPU でデバッグしすくするためとか)
• Thrust とかあるが, でかすぎだったりコンパイル設定しんどい
• std::vector とか std::string とかがぱぱっと使えるくらいでよい
• libcu++(libc++ ベース) https://github.com/NVIDIA/libcudacxx の動きがあるが, 現状 atomic とかだけで vector などコンテナ系は無い
• 2,3 年後くらいには vector とかコンテナ系対応しているかも

NanoSTL をいくつか CUDA 対応してみました
現状 master に取り込んでいます.

移植手順

基本 __CUDACC__ が定義されていたら, 関数に __device__, __host__ 追加するだけでした.

NanoSTL はヘッダだけでライブラリ部分はありませんので, CUDA プログラム側では include するだけで使えます.

サンプル

# include <nanovector.h>

extern "C" {
__global__ void myfunction() {
  nanostl::vector<float> ab;
  ab.push_back(1.0f);
}

という感じに普通に使えます. CUDA(NVCC)自体は, C++ 構文自体は C++17 まで対応していますので, 必要なら using で std namespace にマッピングしてあげて.

C++ モード(pycuda の場合は no_extern_c=1)でコンパイルする必要あります.

制約

host と device でメモリをよろしくやり取りするのは考慮していません.
(e.g. CPU 側で定義した vector を GPU 側から透過的に見えるようにするなど)

明示的にメモリ転送のロジックを記述していただく必要があります.

Note

CUDA では, new/delete, malloc()/free() が device 側から呼べます(正確には SM 2.x(Fermi) 以降で対応であるが, Fermi サポートは打ち切られ現状は Kepler 以上という状況なので, 実質対応していると考えてよい)

OpenCL for C++(OpenCL 3.0)でも, 同様に kernel 側で new/delete or malloc()/free() 相当を呼ぶことができるのであれば比較的簡単に NanoSTL を OpenCL に移植できるでしょう.