CUDA Academy · レッスン

ページ可能メモリが遅い理由

通常のmallocの背後でバッファがステージングされる仕組みを学びます。

レッスン 1/413 ステップ

「ページ可能メモリが遅い理由」は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フィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

  1. ページ可能メモリが遅い理由
  2. cudaMallocHostでピン留めメモリを使う
  3. ストリーム内でcudaMemcpyAsyncを使う
  4. ダブルバッファリングのパイプライン
← CUDA Academyに戻る