CUDAのカーネル内における過剰なメモリ確保について

Question

前提

CUDAを用いてMCTSにおけるシミュレーションの並列化処理を目指しています。MCTSを実装する対象としてBridgetというマイナーなボードゲームを用いています。

実現したいこと

• ゲームの合法手探索によって得られた合法手を格納する配列を長さ3000にて宣言したい。
• 宣言の際に、メモリの確保が過剰に行われているのを確認できたため、その過剰に確保される原因が知りたい。

発生している問題・エラーメッセージ

• 合法手を格納する配列laを宣言する際に確保されるメモリがGPUの持つメモリ量10018[MiB]より大きくなってしまうため、メモリ外アクセスのエラーを吐いてしまう。
• 配列laは大きさ3000のAction_cu型の配列であるため、Action_cuの大きさが56であることを考えると、単純計算で3000 * 56 = 168000[B]のメモリ使用で済むはずである。
• 実際に確保されているメモリ量が膨大になっていることが考えられる。

以下コンパイル後の実行にて得られた出力結果が以下になります。メモリ外アクセスを起こしているとのエラーメッセージを確認できます。

$ nvcc -O2 -lcurand -I./header -g -G  -o obj/kernel.o -c src/kernel.cu
$ nvcc -O2 -lcurand -I./header obj/state.o obj/action.o obj/piece.o obj/measuring.o obj/board.o obj/ucb1.o obj/node.o obj/game.o obj/mcts.o obj/cell.o obj/argmax.o obj/main.o obj/trans_data.o obj/kernel.o obj/playout_cuda.o -o bin/bridget

$ ./bin/bridget
Error: src/playout_cuda.cu:44, code:2, reason: out of memory

該当のソースコード

typedef struct _Piece_cu {
    int kind;
    int rdx;
    int rdy;

    int rcs[3][3];
} Piece_cu;

typedef struct _Action_cu {
    // どのような一手であるのかを示す構造体
    // sizeof(Action_cu)にて大きさを測ったところ56であった。
    int x, y;

    Piece_cu piece;
} Action_cu;

__device__ int playout_gpu(State_cuda *sc, curandState *s) {
    while (1) {
        *** 割愛 ***

        Action_cu la[3000]; // 合法手を格納したい配列を宣言している。

        *** 割愛 ****
    }
}

試したこと

• 配列laのメモリ使用量が膨大であることが原因であるため、少しずつ大きくすることで実際にどれくらいの使用量となっているのか計測することが可能なのではないかと考えた。
• Action_cu la[X]のXを少しずつ大きくしていき、実際にメモリ使用量がどのように変化したのかnvidia-smiを用いて計測した。

配列の大きさX	1	10	100	1000
メモリ使用量[MiB]	339	389	893	5193

• これによって分かるのが、配列の大きさ1あたりで確保されるメモリサイズが5.6[MiB]となっていることである。
• 5.6[MiB]は本来の56[B]と比較すると約10^6倍もの差があることから明らかに不自然な挙動となっている。

バージョン等について

• OS: Ubuntu 20.04.2 LTS
• CUDA Version: 11.3
• GPU: GeForce RTX3080

最後に

こちらのサイトでの質問が初めてなことから書くべき情報が足りていないかもしれません。その際、必要な情報を追記させていただきますので、連絡いただけると幸いです。回答の程よろしくお願いします。

Comments

[matukeso]

総スレッド数は幾つ？
1スレッドあたり158KBで、1e6個thraadを発生させれば158GB必要になる

[koume80703]

返信ありがとうございます。
並列数のことでしたら、デバッグの段階ですので1並列にて実行しております。

[matukeso]

playout_gpu<<<1,1,0,stream>>>(sc,s)ってこと？

[koume80703]

dim3 block(threads_per_block, 1, 1);
dim3 grid((n_elem + block.x - 1) / block.x, 1, 1);
kernel<<<grid, block>>>(省略);

という形で実行しております。threads_per_block, n_elemともに1として実行しています。kernelはホスト側から最初に呼び出す関数として定義しており、そのkernel内からplayout_gpu関数を呼び出しております。
情報記載不足で申し訳ありません。

[jbpb0]

> 1並列にて実行

それを2, 4と増やしたら、配列の大きさ1あたりで確保されるメモリサイズも2倍、4倍になるのでしょうか？

Answer 1

CUDAのSTACK（レジスタからあふれた分）は、GlobalMemoryに確保されますが、いくつ確保されるかというと、「そのデバイスの最大スレッド数」固定です。
たとえば、RTX3080だと、multiProcessorCount=68で、maxThreadsPerMultiProcessor=1536だと、102912個のスタックフレームが個確保されます。

Comments

[koume80703]

私自身の理解が浅く申し訳ないのですが、スタックとは今回確保したい配列laのことでそれが最大スレッド数分だけ確保されるためにメモリを多く食ってしまっているという解釈でよろしいのでしょうか？意図した確保されるべきメモリ量に対して約10万倍ものメモリが確保されていることからmatukeso様のおっしゃる102912個のスタックフレームの確保が影響しているのではないかと思います。

[matukeso]

laが小さければレジスタに載りますが、今のサイズだとその通りですね。
どうしてもというのであれば、外部からスレッド数に応じたlaを与える手です。

[koume80703]

laのサイズは3000の長さを想定しており、その大きさは168000[B]なのですが、レジスタとはそれほどまでに小さいものなのでしょうか？GPUの仕様書を見ればそのレジスタの大きさはわかりますでしょうか？

[matukeso]

deviceQueryすればわかりますが、最近のGPUではMultiProcessorあたり65536個、レジスタはあります(256KB)。これを実行中のスレッドが共有します。スレッドが少ない時でも、1スレッドあたりではmax255個とかです。
ふつうはたくさんスレッドを生やすので、30個とか、おおく使って128個とかになるでしょう。
shared memoryもMultiProcessorあたり48KB〜64KB程度ありますが。

laならまあ、スレッドあたりなんとか2、3個載せれるくらいですね。