ページ可能メモリが遅い理由
通常のmallocの背後でバッファがステージングされる仕組みを学びます。
「ページ可能メモリが遅い理由」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン1/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
データはどこにあるのか
malloc で確保した通常の C++ バッファは、ページング可能なホストメモリ上に置かれます。一見便利ですが、GPU はそこから直接コピーできず、この小さな違いが速度の低下につながります。
ページング可能の意味
メモリがページング可能であるとは、オペレーティングシステムがそのページをいつでも移動したり、ディスクにスワップしたりできるという意味です。柔軟性には優れていますが、固定アドレスを必要とするハードウェアには扱いにくい性質です。
GPU は DMA を使う
GPU はDMAを使ってデータを取得します。これは、移動しない物理アドレスを必要とするハードウェアによる直接転送です。ページング可能なページはこの条件を満たせないため、直接使用できません。
隠れたステージング処理
これを回避するため、ドライバーはまずページング可能なバッファを隠れたステージングバッファにコピーし、その後で GPU に送ります。自分では記述していない余分なコピーのコストを支払うことになります。😮
1 回ではなく 2 回のコピー
つまり、ページング可能なメモリからの 1 回の cudaMemcpy は、実際には2 回のコピーです。ホストからステージングへ、続いてステージングからデバイスへコピーします。この余分な処理こそが、ページング可能な転送が遅く感じられる理由です。
通常の malloc
これは、まず誰もが使おうとするバッファです。動作はしますが、h_data からの転送はすべて、ひそかにステージング処理を経由します。
float* h_data = (float*)malloc(N * sizeof(float));
cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice);OS が気にする理由
OS は、多数のプログラムを同時に動かせるよう RAM をオーバーコミットするため、メモリをページング可能な状態に保ちます。この利便性が、GPU によるページの直接読み取りを妨げています。
失われる帯域幅
ページング可能な転送では、リンクのピーク帯域幅の半分程度しか出ないことがよくあります。ステージングコピーによって CPU サイクルとメモリ帯域が消費され、本来の転送に使えるリソースが減るためです。
重ね合わせも妨げる
ステージングコピーは同期的に行われるため、ページング可能な転送をカーネルと本当に重ね合わせることはできません。ストリームを有効に活用する非同期処理の利点が失われます。
それでも問題になる場面
小さなコピーを 1 回行うだけなら、誰も気にしないでしょう。しかし、毎回のループでデータを転送する場合は、ステージングのコストが積み重なり、実行時間の大部分をひそかに占めるようになります。
修正がやって来ます
解決策は、ページアウトできないバッファー、つまりピン留めメモリを CUDA に要求することです。GPU はステージングも二重コピーもせず、直接読み取れます。
簡単な確認
通常のページング可能なメモリからの転送が、必要以上に遅くなるのはなぜですか?
振り返り
ページング可能なバッファーは移動する可能性があるため、GPU は DMA で直接アクセスできません。ドライバーは最初にステージングするため、コピーが二重になります。この先の解決策がピン留めメモリです。🚀
AI チューターと学ぶ C++ — 無料
ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。
- コース
- 30
- レッスン
- 120
よくある質問
「ページ可能メモリが遅い理由」レッスンは無料ですか?
はい。「ページ可能メモリが遅い理由」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し(組み込みコードエディタと24時間対応のAIチューター)、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「ページ可能メモリが遅い理由」で何を学びますか?
通常のmallocの背後でバッファがステージングされる仕組みを学びます。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか?
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン1/4です。
「ページ可能メモリが遅い理由」レッスンにはどのくらい時間がかかりますか?
ほとんどのCoddyKitレッスンは約5~10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか?
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。