トップページ

最新記事

CUDAのカーネル関数内mallocの挙動について

なんの話か

先日「Fortran/OpenACCのループ内サブルーチンの可変長配列の実体について」という記事を書き、OpenACC内の可変長配列がmallocによって領域確保されている話をしました。
では、これがmallocだとなにか悪いことが起きるの？mallocだし遅かったりするの？という話です。
結論から言うと、スレッド数が多くなるに連れ遅くなります。

malloc/freeにかかる時間の調査

CUDA/C++で1スレッドがfloat用の領域を1つだけmallocしfreeするだけのカーネル関数を書き、それぞれの関数の実行にかかったclock数を計算します。
Grid sizeを1に固定し、Block sizeを1から1024に1ずつ変化させていった場合のclock数の変化を調べます。
比較対象として、同じGrid/Block sizeで立ち上げるが、threadIdx.x==0のスレッドだけがBlock size個分のfloatを1回mallocとfreeするカーネル関数を作成し、同様にclock数を計算します。
つまり、両カーネル関数はトータルで同じ量のメモリをmallocしfreeすることとなります。
ソースコードはこちらです。enp1s0/cuda-in-kernel-malloc - GitHub

で、結果（各スレッドでのclock数の平均）を図示したのがこちらです。
縦軸は4 Byteをmalloc/freeするのにかかる平均clock数ということとなります。

▲ 図1: 各スレッドがmallocを呼ぶ場合と、1スレッドだけが同量の領域を確保するmallocを1回呼ぶ場合の経過clock

トータルで同じサイズ分のmallocをしていますが、やはり全スレッドがそれぞれmallocした場合はスレッド数に比例してclock数がかかっていることがわかります。
おそらくですが、malloc内の一部でatomicが呼ばれて、シーケンシャルに処理をしないといけないところがあるのかなという気持ちです。
そもそもSMからどうやってメモリ確保をしているのかが分からないので、なんとも言えませんね。
もしかしたらenvytoolsのdocumentのどこかに書いてあるかも。

比較対象として1スレッドだけがmallocを呼ぶカーネル関数を用いていますが、これは1回のmallocでBlock size個分のfloat領域を確保します。
これを1回あたりfloat 1個分の確保を行うmallocをBlock size回呼ぶよう変更したら、計算時間はどうなるのでしょうか？
線形に増えることは予想できますが。
というわけで調べてみたのがこちらです。

▲ 図2: 各スレッドが4 byteのmallocを呼ぶ場合と、1スレッドだけが4 byteを確保するmallocをBlock size回呼ぶ場合の経過clock

傾きが小さく見づらいですが、赤線と緑線において、Block sizeが大きくなるに連れclock数も大きくなっていることがわかります。
とはいえ全スレッドがmallocを呼ぶ場合と比較するとそのclock数が小さいです。

おわりに

全スレッドがmallocを呼ぶのは遅いということになりますが、どうしても使わないといけない場合もあるかと思います。
そういうときは同一Warp内でatomicを呼ばないようにするなど、工夫の使用はあるかもしれないです。

記事作成日:2021-09-26