CUDA Academy · レッスン

大きな画像向けにタイルをストリーミングする

I/O と計算をオーバーラップさせる

レッスン 3/413 ステップ

「大きな画像向けにタイルをストリーミングする」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン3/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。

画像が大きすぎる場合

巨大な画像はGPUメモリに収まらなかったり、画像全体をコピーするとパイプラインが停止したりすることがあります。その場合は画像をタイルに分割し、順番に処理します。🖼️

チャンクに分割する

画像を行方向の帯、または矩形のタイルに分割します。各タイルは、デバイスメモリに過度な負荷をかけずに、アップロード、処理、ダウンロードできる大きさにします。

タイルごとの3つの手順

すべてのタイルは同じ手順で処理します。デバイスにコピーし、カーネルを実行し、結果をコピーして戻します。画像全体の処理が終わるまで繰り返します。

タイルを直列処理すると時間を無駄にする

タイルを1つずつ処理すると、コピーのたびにGPUがアイドル状態になります。性能を高めるには、あるタイルの転送と別のタイルの計算をオーバーラップさせる必要があります。

ストリームでオーバーラップを実現する

各タイルの処理を、それぞれ別のストリームに投入します。これにより、あるタイルのコピーがまだ実行中でも、ドライバーは別のタイルのカーネルを実行できます。

cudaStream_t s;
cudaStreamCreate(&s);

非同期コピーが必要

通常のcudaMemcpyはホストをブロックし、オーバーラップを妨げます。ストリーム上でcudaMemcpyAsyncを使用し、すべての処理を停止せずにコピーをキューへ投入してください。

cudaMemcpyAsync(d_tile, h_tile, bytes, cudaMemcpyHostToDevice, s);

ホストバッファをピン留めする

非同期転送でDMAを使用するには、ホストメモリをピン留めする必要があります。cudaMallocHostでタイル用のステージングバッファを確保しないと、オーバーラップが気づかないうちにブロッキング処理へフォールバックします。

cudaMallocHost(&h_tile, bytes);

パイプラインをダブルバッファリングする

2つのタイル用バッファを用意し、ストリームを交互に使います。タイルNを計算している間にタイルN+1をアップロードし、処理の裏で転送を隠す滑らかなパイプラインを構築します。

タイルの境界にハローピクセルを追加する

タイルの境界付近でぼかしを行うには、隣のタイルのピクセルが必要です。重なり部分となるハロー領域を含め、境界付近の結果が正しくなるようにします。

保存前に同期する

非同期処理は、呼び出しが戻った時点ではまだ完了していません。ホスト側でタイルの結果を読み戻す前に、各ストリームに対してcudaStreamSynchronizeを呼び出してください。

cudaStreamSynchronize(s);

大きな画像でもGPUを安定稼働させる

ストリームとダブルバッファリングを使ってタイルを処理すれば、画像のサイズに関係なくGPUを処理で満たし続けられます。転送が計算処理の裏に隠れ、スループットも高く保たれます。⚡

確認問題

タイルを別々のストリームで処理していますが、オーバーラップが見られません。最も可能性の高いミスは何でしょうか。

まとめ

大きな画像をタイルに分割し、非同期コピーとピン留めメモリを使ったストリーム処理でオーバーラップさせ、パイプラインをダブルバッファリングし、境界を正しく処理するためにハロー領域を追加する方法を学びました。🎯

無料で開始

AI チューターと学ぶ C++ — 無料

ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。

コース
30
レッスン
120

よくある質問

「大きな画像向けにタイルをストリーミングする」レッスンは無料ですか?

はい。「大きな画像向けにタイルをストリーミングする」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し(組み込みコードエディタと24時間対応のAIチューター)、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。

「大きな画像向けにタイルをストリーミングする」で何を学びますか?

I/O と計算をオーバーラップさせる ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。

CUDA Academyを始めるのに経験は必要ですか?

事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン3/4です。

「大きな画像向けにタイルをストリーミングする」レッスンにはどのくらい時間がかかりますか?

ほとんどのCoddyKitレッスンは約5~10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。

このCUDA Academyレッスンでコードを書いて実行できますか?

はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。

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

  1. 処理パイプラインを設計する
  2. フィルターを 1 つのカーネルに融合する
  3. 大きな画像向けにタイルをストリーミングする
  4. プロファイル、最適化、出荷
← CUDA Academyに戻る