R&D チームの奥村（@izariuo440）です。今年も新卒が CUDA プログラミングを嗜んでいたところ、思ったより速度が出ないという話を聞いたので「CUDA で一時領域の確保・破棄を回避して速度低下を防ぐ」という対策をまとめました。Happy CUDing!

tech-blog.optim.co.jp tech-blog.optim.co.jp

背景と動機

CUDA でソートや数え上げなどに代表される reduction 操作を実行するとき、一時領域を要する場合があり、そのために一次領域の確保・破棄が必要になることがあります。繰り返し一時領域の確保・破棄が呼び出されると、数ミリ秒程度の実行コストがかかることがあります。例えば、全体の処理時間が 16 ミリ秒かかっているとして、ここに 1 ミリ秒の実行コストが加わると、FPS は 62.5 から 58.9 に落ちます。4 ミリ秒だと、50 fps まで落ちてしまいます。これをうまく回避して、速度低下を防ぎます。

CUDA の同期について

CUDA の同期には、「明示的な同期」と「暗黙的な同期」があります。

明示的な同期 (Explicit Synchronization¹)
- プログラマが同期を行うコードを書いて実行される同期
  1. . cudaDeviceSynchronize()
  2. . cudaStreamSynchronize()
  3. . cudaStreamWaitEvent()
  4. . etc...

暗黙的な同期 (Implicitly Synchronization²)

ホストスレッドから、以下の CUDA 命令³を実行したときに実行されてしまう同期

1. Page-locked memory allocation
     1. cudaMallocHost
     1. cudaHostAlloc
1. Device memory allocation
     1. cudaMalloc
1. Non-Async version of memory operations
     1. cudaMemcpy* (no Async suffix)
     1. cudaMemset* (no Async suffix)
1. Change to L1/shared memory configuration
     1. cudaDeviceSetCacheConfig

暗黙的な同期について理解しておくと「なんか思ったよりも速度が出ないな」という状態に陥りがちです。そうしたときは、上記が含まれていないか疑ってみるとよいでしょう。

暗黙的な同期を防ぐ基本方針は「メモリ確保・破棄を避ける」と考えてもよいと思います。element-wise な CUDA カーネルを書く程度なら、一度確保したメモリを使い回すなどして、プログラマが気をつければ事足ります。

暗黙的なメモリ確保・破棄を回避する

注意が必要なのは、thrust⁴ などのライブラリを使うときです。

例えばソート。自分で実装するのは大変⁵なので、ライブラリにある関数を使うことが多いと思います。例えば thrust だと、ソートは基数ソートで実装されていますが、これは一時領域を必要とします。一時的な領域のためのメモリ確保・破棄を行うため、暗黙的な同期が実行されます。また、外部から一時領域を指定することはできません。

しかし、thrust はよくできていて、一時領域の確保・破棄にうまく割り込むことができます。C++ の標準ライブラリのように、カスタムアロケータを仕込むことができるのです。thrust の各種アルゴリズムは第一引数に実行ポリシーを指定するオーバーロードを持っており、ここにカスタムアロケータを関連付けることができるのです。以下にコードのイメージを記載します。

// 1. CUDA でデフォルトストリームで実行するポリシー
auto exec1 = thrust::cuda::par;

// 2. CUDA で指定したストリームで実行するポリシー
auto exec2 = thrust::cuda::par.on(stream);

// 3. CUDA で指定したアロケータを使うポリシー
auto exec3 = thrust::cuda::par(alloc);

// 4. CUDA で指定したアロケータを使い指定したストリームで実行するポリシー
auto exec4 = thrust::cuda::par(alloc).on(stream);

上記の 3 と 4 のようにすると、カスタムアロケータを使うことができます。

Thrust/CUDA tip: reuse temporary buffer across multiple transforms には、カスタムアロケータを指定した効果がコードつきで解説されています。以下に意訳しておきます。

Thrust は GPU で高速にデータ処理するための便利な STL っぽいテンプレートライブラリだ。実際のアプリケーションでは、同じデータ処理を複数回異なるデータセットで実行することがよくある。reduction を含む transform は一時領域の確保が必要だ。既定では確保は transofrm ごとに実行されて、cudaMalloc/cudaFree が呼び出され、結構コストが高い。

複数の transform にまたがって一つの一時領域を使い回すのが明らかにいい方法だろう。これは cudaMalloc/cudaFree をアプリケーションレベルでフックすればできる。Thrust の場合、もっといい解決策がある。カスタムアロケータを使うんだ。

コードは省略。

このカスタムアロケータを使うと、cudaMalloc/cudaFree が一度しか呼び出されなくなる。

我々が開発した金融アプリでは、この手法で 8% 性能を改善した。

というわけで、CUDA の速度改善あるあるといえるのではないでしょうか。

thrust のサンプルコードにもカスタムアロケータの例があります。

カスタムアロケータ

カスタムアロケータは、以下のように実装できます。

class my_allocator {
public:
  typedef char value_type;

  char* allocate(std::ptrdiff_t num_bytes) {
    // TODO: implement.
  }

  void deallocate(char* ptr, size_t n) {
    // TODO: implement.
  }
};

allocate() に cudaMalloc()、deallocate() に cudaFree() を使うことでカスタムアロケータとして動作させることができますが、それだとメモリ確保・破棄を防ぐことができません。

私は、予め大きめのメモリ領域を確保して（メモリプール）、そこから小さい領域に区切って利用する方式⁶をとることで、繰り返しメモリが確保・破棄されることを防いだりします。動画のリアルタイム解析などでは、毎フレームごとにリセットする機会があるので、そこでポインタを巻き戻しておきます。

メモリプールを超えるような確保要求があった場合は、やむなく動的にメモリを確保します。動的にメモリを確保した分は、リセットするタイミングでメモリプールをその分だけ大きくして、次回から動的なメモリの確保が発生しないようにしています。

メモリアロケータの世界は奥が深いです。ゲーム界隈でもこうした取り組みが行われているので、興味のある方は以下を参照すると面白いと思います。

カスタムメモリマネージャと高速なメモリアロケータについて

Buddy memory allocation なんかも有名だと思いますが、実装が面倒なので私は使いませんでした。

Buddy memory allocation - Wikipedia

おわりに

この記事は、数年前に社内のナレッジベースに書いていたものがベースですが、カスタムアロケータについて加筆修正しました。

オプティムでは、こうした技術に興味がある・作ってみたい・既に作っている、というエンジニアを募集しています。興味のある方は、こちらをご覧ください。

www.optim.co.jp

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#explicit-synchronization ↩
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#implicit-synchronization ↩
CUDA C/C++ Streams and Concurrency ↩
NVIDIA が公開している CUDA 対応の STL っぽい並列計算テンプレートライブラリ↩
CPU でソートするのと GPU で高速にソートするのはまったく別世界です。↩
小さい領域に区切るとき、アラインメントが１６バイト境界になるよう注意する必要があります。↩