著者 |ルーカス・デ・リマ・ノゲイラ 編纂者：岳陽 AI の支援を受けて著者が作成した画像 (https://copilot.microsoft.com/images/create) 今日では、ディープラーニングというと、パフォーマンスを向上させるために GPU を使用することを考えるのが自然です。 GPU（グラフィックス・プロセッシング・ユニット）は、もともと画像や2D・3Dグラフィックスのレンダリングを高速化するために開発されました。しかし、その強力な並列計算能力により、GPUの用途はディープラーニングなどの分野へと急速に拡大しています。 ディープラーニングモデルへのGPUの応用は2000年代半ばから後半にかけて始まり、2012年のAlexNetの登場がこのトレンドをさらに加速させました。Alex Krizhevsky氏、Ilya Sutskever氏、Geoffrey Hinton氏によって設計・開発された畳み込みニューラルネットワークであるAlexNetは、2012年のImageNet Large Scale Visual Recognition Challenge (ILSVRC)で華々しいデビューを飾りました。この勝利は画期的な出来事であり、画像分類におけるディープニューラルネットワークの優れた性能を実証しただけでなく、大規模モデルの学習にGPUを使用することの有効性を浮き彫りにしました。 この技術的進歩に続いて、GPU はディープラーニング モデルでますます広く使用されるようになり、PyTorch や TensorFlow などのフレームワークが登場しました。 現在、PyTorchで「.to("cuda")」と入力するだけでデータをGPUに渡し、モデルの学習を高速化できます。しかし実際には、ディープラーニングアルゴリズムはGPUのコンピューティングパワーをどのように巧みに活用しているのでしょうか？早速見ていきましょう。 ニューラルネットワーク、CNN、RNN、Transformerといったディープラーニングのコアアーキテクチャは、行列の加算、乗算、行列への関数の適用といった基本的な数学演算を基盤としています。したがって、これらのコア演算を最適化することが、ディープラーニングモデルのパフォーマンス向上の鍵となります。 最も基本的なシナリオから始めましょう。2つのベクトルに対して加算演算 C = A + B を実行する必要があるとします。 この機能は、C プログラミング言語を使用して簡単に実装できます。 従来、コンピュータはベクトルの各要素に一つずつアクセスし、各反復処理で各要素のペアを順番に加算する必要があることは容易に理解できます。しかし、要素のペア間の加算演算は互いに独立していることに注意することが重要です。つまり、ある要素のペアの加算は、他の要素のペアに依存しません。では、これらの数学演算を同時に実行し、すべての要素のペアの並列加算を実現できたらどうなるでしょうか？ 最も単純なアプローチは、CPUのマルチスレッド機能を活用してすべての数学演算を並列に実行することです。しかし、ディープラーニングでは、数百万要素にも及ぶ巨大なベクトルを処理する必要があります。通常、一般的なCPUは12個程度のスレッドを同時に処理することしかできません。ここでGPUの優位性が明らかになります。現在主流のGPUは数百万のスレッドを同時に実行できるため、大規模ベクトルにおける数学演算の効率を大幅に向上させることができます。 01 GPUとCPUの比較単一操作の処理速度ではCPUがGPUにわずかに優位に立つかもしれませんが、GPUの優位性は優れた並列処理能力にあります。この違いは、両者の設計上の優先順位の根本的な違いに起因しています。CPUは単一の操作シーケンス（つまりスレッド）を効率的に実行するように設計されており、一度に処理できるのは数十個程度です。一方、GPUは数百万個のスレッドを並列処理するように設計されており、個々のスレッドの速度を犠牲にして、全体的な並列性能を飛躍的に向上させています。 例えば、CPUをクールなフェラーリのスポーツカー、GPUを広々としたバスに例えてみましょう。乗客1人を輸送するだけのタスクであれば、フェラーリ（CPU）が間違いなく最適な選択肢です。しかし、もし複数の乗客を輸送する必要がある場合、片道の速度はフェラーリ（CPU）の方が速いものの、バス（GPU）は一度にすべての乗客を収容できるため、バス全体の輸送効率はフェラーリのように複数の乗客を個別に輸送する効率をはるかに上回ります。つまり、 CPUは連続した単一のタスクを処理するのに適しており、GPUは多数のタスクを並列処理する際に優れた性能を発揮します。 AI の支援を受けて著者が作成した画像 (https://copilot.microsoft.com/images/create) 優れた並列コンピューティング機能を実現するために、GPUはデータキャッシュやフロー制御メカニズムよりもデータ処理に多くのトランジスタを割り当てるように設計されています。これはCPUとは全く異なる設計思想です。CPUは、単一スレッドの実行効率を最適化し、複雑な命令セットを処理するために、意図的に多数のトランジスタを割り当て、これらの領域のパフォーマンスを向上させます。 下の図は、CPU と GPU 間のチップ リソース割り当ての大きな違いを鮮明に示しています。 画像はCUDA C++プログラミングガイドからインスピレーションを得た著者によるものです。 (https://docs.nvidia.com/cuda/pdf/CUDA_C_プログラミングガイド.pdf) CPUは強力なコアと、より洗練されたキャッシュメモリアーキテクチャ（大量のトランジスタリソースを消費）を備えています。この設計により、シーケンシャルタスクの実行速度が大幅に最適化されます。一方、グラフィックスプロセッサ（GPU）は、コア数を重視し、より高い並列処理能力を実現します。 ここまで基礎を解説してきましたが、実際のアプリケーションで並列コンピューティングの利点を効果的に活用するにはどうすればよいでしょうか。 02 CUDA入門ディープラーニングモデルを構築する際には、PyTorchやTensorFlowといった人気のPythonライブラリが選ばれることが多いでしょう。しかし、これらのライブラリのコアコードはC/C++で記述されていることは否定できません。さらに、前述のように、GPUを活用してデータ処理を高速化することは、多くの場合、最適化戦略の主流となっています。ここでCUDAの重要性が明らかになります。CUDAはCompute Unified Device Architectureの略で、NVIDIAが汎用コンピューティングにおいてGPUを最大限に活用できるように綿密に構築したプラットフォームです。ゲームエンジンでグラフィカルコンピューティングに使用されるDirectXとは異なり、CUDAは開発者がNVIDIAのGPUコンピューティングパワーを、グラフィックスレンダリングだけでなく、汎用ソフトウェアに統合することを可能にします。 この目標を達成するために、CUDA は、開発者が GPU 仮想命令セットを呼び出して特定の操作 (CPU と GPU 間でのデータ転送など) を実行できるように、シンプルな C/C++ ベースのインターフェイス (CUDA C/C++) を導入しました。 技術的な詳細に入る前に、CUDA プログラミングにおけるいくつかの基本的な概念と技術用語を明確にする必要があります。
したがって、CUDAを使用して記述された基本コードでは、メインプログラムはホスト（CPU）上で実行され、デバイス（GPU）にデータを渡し、デバイス（GPU）上で並列実行されるカーネル（関数）を呼び出します。これらのカーネルは複数のスレッドによって同時に実行されます。計算が完了すると、結果はデバイス（GPU）からホスト（CPU）に返されます。 とはいえ、2 つのベクトルのセットを追加するという具体的な問題に再び焦点を当ててみましょう。 CUDA C/C++ を使用すると、プログラマーはカーネルと呼ばれる C/C++ 関数を作成できます。これらのカーネルが呼び出されると、N 個の異なる CUDA スレッドがそれらを N 回並列に実行します。 このタイプのカーネルを定義するには、宣言指定子として 各CUDAスレッドには、カーネル実行時に一意のスレッドID「threadIdx」が割り当てられます。このIDは、カーネル内の定義済み変数から取得できます。上記のサンプルコードは、長さNの2つのベクトルAとBを加算し、その結果をベクトルCに格納します。ループ内でペアごとの加算を順番に処理する従来の逐次的なアプローチと比較して、CUDAの利点は、N個のスレッドを並列に利用してすべての加算処理を一度に完了できることにあります。 ただし、上記のコードを実行する前に、一度変更する必要があります。カーネル関数はデバイス（GPU）上で実行されるため、関連するすべてのデータはデバイスのメモリに格納されている必要があることに注意してください。これを実現するために、CUDAが提供する以下の組み込み関数を使用できます。 変数A、B、Cをカーネルに直接渡すことは、この状況には適していません。ポインタを使用する必要があります。CUDAプログラミング環境では、ホスト配列（例のA、B、Cなど）をカーネルの起動に直接使用することはできません（<<<…>>>）。CUDAカーネルのワークスペースはデバイスメモリであるため、デバイスメモリ上でカーネルが確実に実行されるようにするには、デバイスポインタ（d_A、d_B、d_C）をカーネルに提供する必要があります。 さらに、cudaMalloc 関数を呼び出してデバイス上のメモリ領域を割り当て、cudaMemcpy を使用してホストとデバイス間のデータ転送を実装する必要があります。 これで、コード内でベクトル A と B を初期化し、プログラムの最後に CUDA メモリをクリーンアップできるようになりました。 さらに、カーネルを呼び出した後、必ず さらに、 CUDAのエラー検出メカニズムも同様に不可欠です。このメカニズムは、GPU上の潜在的なプログラム欠陥（バグ）を迅速に特定し、修正するのに役立ちます。このステップを無視すると、デバイススレッド（CPU）は実行を継続し、CUDA関連の問題のトラブルシューティングが非常に困難になり、CUDA関連のエラーを特定することが困難になります。 これら 2 つのテクノロジの具体的な実装方法は次のとおりです。 CUDAコードをコンパイルして実行するには、まずCUDAツールキットがシステムにインストールされていることを確認する必要があります。次に、NVIDIA CUDAコンパイラであるnvccを使用して、必要なコードをコンパイルします。 しかし、現在のコードにはまだ最適化の余地があります。前の例では、N = 1000のベクトルを処理しましたが、これはGPUの強力な並列処理能力を十分に発揮するには小さすぎます。特にディープラーニングのシナリオでは、数百万のパラメータを含む大規模なベクトルを扱うことがよくあります。しかし、Nの値を500,000に設定し、<<<1, 500,000>>>形式でカーネルを実行しようとすると、上記のコードはエラーをスローします。したがって、コードを改善し、このような大規模な操作を正常に実行できるようにするには、CUDAプログラミングの中核概念であるスレッド階層を早急に習得する必要があります。 03 スレッド階層カーネル関数を呼び出す際の表記は、<<<ブロック数, ブロックあたりのスレッド数>>>という形式です。したがって、上記の例では、N個のCUDAスレッドを単一のスレッドブロックとして起動しました。ただし、同じスレッドブロック内のすべてのスレッドが同じストリーミングマルチプロセッサコア上に共存し、そのメモリリソースを共有する必要があるため、各スレッドブロックが保持できるスレッド数には制限があります。 この制限の特定の値を照会するには、次のコードを使用できます。 著者が現在使用しているGPUは、スレッドブロックあたり最大1024スレッドしかサポートできません。そのため、例に示したような巨大なベクトルを効率的に処理するには、より多くのスレッドブロックを配置し、より大規模な同時スレッド実行を実現する必要があります。これらのスレッドブロックは、次の図に示すように、グリッド状に慎重に配置されています。 https://handwiki.org/wiki/index.php?curid=1157670 (CC BY-SA 3.0) 以下の方法でスレッド ID を取得できるようになりました。 したがって、コード スクリプトは次のように更新されました。 04 パフォーマンス比較分析以下の表は、異なるサイズのベクトルに対して加算演算を実行する場合の CPU と GPU の計算パフォーマンスの比較を示しています。 著者による画像 GPUの処理能力の優位性は、大規模なベクトルを扱う場合にのみ顕著になることは明らかです。さらに、この時間比較はカーネル/関数の実行時間のみを考慮しており、ホストとデバイス間のデータ転送に必要な時間は考慮していないことを理解することが重要です。データ転送時間はほとんどの場合無視できますが、単純な加算演算のみを実行する現在のシナリオでは、比較的大きな影響を及ぼします。したがって、 GPUの計算性能は、計算負荷が高く、かつ超並列処理に適したタスクに直面した場合にのみ、真に発揮されることを忘れてはなりません。 05 多次元の糸これで、単純な配列演算のパフォーマンスを向上させる方法がわかりました。しかし、ディープラーニングモデルを扱う場合、行列演算とテンソル演算も処理する必要があります。これまでの例では、N個のスレッドを含む1次元ブロックのみを使用しました。しかし、多次元スレッドブロック（最大3次元をサポート）の実行も同様に可能です。したがって、便宜上、行列演算を処理する必要がある場合は、N x M個のスレッドで構成されるスレッドブロックを実行できます。行列の行インデックスは `row = threadIdx.x` を使用して決定でき、列インデックスは `col = threadIdx.y` を使用して取得できます。さらに、演算を簡素化するために、`number_of_blocks` と `threads_per_block` を `dim3` 変数型を使用して定義できます。 以下のサンプル コードは、 2 つの行列の加算演算を実行する方法を示しています。 さらに、この例を拡張して複数のスレッド ブロックを処理することもできます。 さらに、同じアプローチを使用して、この例を 3 次元操作の処理に拡張できます。 前のセクションでは、多次元データの処理方法を紹介しました。次に、重要でありながら理解しやすい概念、つまりカーネル内で関数を呼び出す方法を学びます。これは通常、 CUDAプログラミングのコアコンセプトを理解できたので、CUDAカーネルの構築に取り掛かることができます。ディープラーニングモデルの場合、これは基本的に行列とテンソルに対する一連の演算（加算、乗算、畳み込み、正規化など）を伴います。例えば、基本的な行列乗算アルゴリズムは、次のように並列化できます。 GPU版の行列乗算アルゴリズムではループ回数が大幅に削減され、処理速度が大幅に向上していることがわかります。次のグラフは、CPUとGPUにおけるN x N行列乗算のパフォーマンスを視覚的に比較したものです。 著者による画像 行列のサイズが大きくなるにつれて、行列乗算演算を処理する際の GPU のパフォーマンスがさらに向上することがわかります。 次に、基本的なニューラル ネットワーク モデルに注目しましょう。このモデルの基本操作は通常、次の図に示すように y = σ(Wx + b) と表されます。 著者による画像 上記で説明した演算は、主に行列の乗算、行列の加算、そして配列への関数の適用です。これらの並列処理技術を習得すれば、GPU上でニューラルネットワークをゼロから構築できるようになります。 06 結論この記事では、GPU処理によるディープラーニングモデルのパフォーマンス向上の基本概念について解説します。ただし、この記事はあくまでも表面的な説明に過ぎず、その背後にはより深い技術が隠されていることに留意してください。PyTorchやTensorflowなどのフレームワークは、最適化されたメモリアクセスやバッチ処理（cuBLASやcuDNNなどのCUDAベースのライブラリを利用）といった複雑な概念を含む、数多くの高度なパフォーマンス最適化手法を実装しています。この記事が、読者の皆様に、`.to("cuda")`メソッドを用いてGPU上でディープラーニングモデルを構築・実行する際の基本原理について、基本的な理解を深めていただければ幸いです。 読んでくださってありがとうございます！😊 ルーカス・デ・リマ・ノゲイラ https://www.linkedin.com/in/lucas-de-lima-nogueira/ 終わり オリジナルリンク: https://towardsdatascience.com/why-deep-learning-models-run-faster-on-gpus-a-brief-introduction-to-cuda-programming-035272906d66 |
01 GPUとCPUの比較02 CUDA入門03 スレッド階層04 パフォーマンス比較分析05 多次元の糸06 結論 |
