それで、グラフィックカードのプログラミングに関する最初の投稿と、これがどれほど複雑かについてのホラーストーリーからほぼ1年が経過しました。 ここで、すべてがそれほど悪くないことと、OpenCLと呼ばれるこの奇妙なものを使用する方法を示し、さらにその主な利点、つまり異なるデバイスで同じコードを実行する機能を使用することを示します。 また、従来のプロセッサのパフォーマンスをほぼ無料で1桁向上させる方法も示します。

はじめに

WikipediaでOpenCLについて語るのはあまり意味がないと思いますが、簡単に言えば、OpenCLは異なるアーキテクチャ、特に高度な並列プロセッサで異なるデバイスで同じコードを実行できる言語（フレームワークとプラットフォーム）です、ビデオカードや最新の中央処理装置など。 この標準はC99に基づいており、クロノスグループによってサポートされています。これに基づいて、教育プログラムは完了したと見なします。



OpenCLがどのように機能するかについて話している間に、小さなコードを見せてそこで何が起こっているのかを説明することから始めます。



まず、ささいなコードについて説明します。OpenCLの魔法を見たくない人は、最初の部分をスキップできます（MathCalculations関数について説明する最後の段落を読んでください。これは重要です。 5番目のセクションに直接進みますが、とにかくMathCalculationsを見てください。



これが、OpenCLをテストするための、またはむしろ、後で説明する特定の抽象的な数学式を計算するための私の小さなプログラムの主な外観です。 そこで、ここで何が起こっているのかを理解するために、行ごとに見てみましょう。

パート1-ソースデータと従来のコンピューティング方法の初期化

GenerateTestData（）; 特別なことは何も行いませんが、単に入力および出力配列にメモリを割り当て、入力配列をランダムデータで満たします。



次に、もう少し興味深い機能があります。



その最初のサイクル

 for(int iTest=0; iTest<(TESTS_NUMBER/10); iTest++)
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

より正確なランタイムを得るために、テストを数回実行する必要がありました。 各テストの計算時間はtimeValues配列に格納され、そこから平均値が計算されてhostPerformanceTimeMSに格納されます。



第二サイクル

 for(int iJob=0; iJob<DATA_SIZE; iJob++)
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

入力配列の要素に対していくつかの数学計算を順次実行し、出力配列に格納します。



ご覧のとおり、このコードには異常なものはありません。通常のシステムコンパイラによってコンパイルされ、毎日作成するほとんどのコードのように、中央処理装置で順次実行されます。 そして、OpenCLによって得られた結果を後で彼と比較し、どのようなパフォーマンスの向上が得られるかを理解するために彼が必要です。



すぐにMathCalculationsを調べ、すべてが完全に退屈していることを確認する必要があります。



実際、それはあまり意味がありません（そして、非常に単純化できることは明らかです）が、純粋な数学演算の簡単なデモンストレーションとして機能します。 その中で重要なことは、それが別の.cppファイルにあり、多くの算術演算を行うことですが、それについては後で詳しく説明します。

パート2-OpenCLの初期化

だから、患者はこの部分を読んで、何か面白いことが始まったことを喜んでいたが、せっかちな人はこの感覚を経験することができなかったので、最後の段落をスキップした:)



最初に、OpenCLランタイムAPIはC ++のAPIではなく、CのAPIにすぎないと言います。 一般に、エラーをチェックするために、各関数によって返されたコードをチェックする必要があることを除いて、これには何も問題はありません。これはあまり便利ではありません。 また、割り当てられたリソースのリリースを手動で監視する必要があります。

しかし、公式のC ++ラッパー（KhronosのWebサイトにあります）もあります。これは、OpenCLオブジェクトに対応し、参照カウント（参照）をサポートし、エラーの場合に例外をスローするクラスのセットです（#define __CL_ENABLE_EXCEPTIONSを使用して例外を有効にする必要があります） ） この非常にラッパーは、テストで使用します。



したがって、最初に入手できるのは、利用可能なプラットフォームのリストです。

 vector<cl::Platform> platforms; cl::Platform::get(&platforms);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

OpenCLのプラットフォームはベンダーに対応しています。 NVidiaには1つのプラットフォームにデバイスが、Intelには別のプラットフォームなどがあります 私の場合、利用できるのは2つのNVidiaとIntelプラットフォームだけです。



もう1つの小さなトリックは、C ++ラッパーが独自のベクター（それについて彼に伝える場合）またはSTDのベクターを使用できるため、例でcl :: vectorのようなものが出てきても、恐れないで両方の形式を知っている。



プラットフォームのリストを取得した後、プラットフォームごとに使用可能なデバイスのリストを取得します。

 std::vector<cl::Device> devices; platforms[iPlatform].getDevices(CL_DEVICE_TYPE_ALL, &devices);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

実際、デバイスは計算を実行するものです。 GPU、CPU、またはホストに接続されている特別なアクセラレーター、つまり OpenCLが実行されているシステム。 CL_DEVICE_TYPE_ALLの代わりに、CL_DEVICE_TYPE_GPUを渡すことができます。そうすると、中央処理装置用にビデオカードまたはCL_DEVICE_TYPE_CPUのみが発行されます。



私が見つけた各デバイスについて、以下で説明するテストを実行し、問題が発生した場合にOpenCLがスローする例外をキャッチしようとします。すべてがうまくいった場合、CheckResultsは結果をホストの最初の部分でカウントしたものと比較し、統計を計算します間違い。

パート3-カーネルの作成と起動

ここで最も興味深い部分-計算に進みます。



まず、この方法で取得したデバイス名を表示します。

 device.getInfo<CL_DEVICE_NAME>()
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

同様に、コアの数、周波数、バージョンなどに関する情報を取得できます。



次に、コンテキストを作成します。

 vector<cl::Device> contextDevices; contextDevices.push_back(device); cl::Context context(contextDevices);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

コンテキストではすべてがそれほど単純ではありません...コンテキストを作成するとき、それに含めるデバイスのリストを渡しますが、制限があります：同じプラットフォーム上のデバイスのみが同じコンテキストに入れることができます。 GPUおよびCPU（Intel / NVidiaの場合）とのコンテキストの作成は失敗します。 同じコンテキストに複数のデバイスがある場合、すべてのバッファは異なるデバイスで自動的に同期されます。 一方では、これによりマルチGPUサポートが簡素化され、他方では、ドライバーがどのように、何を、いつ同期するかが誰にもわかりません。データ転送の効率は、すべてが考えられる高いパフォーマンスを得るために重要です。 そのため、通常は各デバイスに個別のコンテキストを作成し、データを手動で配布します。 したがって、何が、どこで、いつ発生するかは常にわかっています。



次のステップは、デバイスのコマンドキューを作成することです。

 cl::CommandQueue queue(context, device);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

このキューは特定のデバイスに関連付けられており、理論的には故障している可能性がありますが、実際にはこの動作に気付きませんでした。 1つのデバイスに複数のキューがあり、同じコンテキスト内で異なるキューからコマンドを同期できます。



次に、入力ベクトルと出力ベクトル用のバッファーを作成します。

 //Create memory buffers cl::Buffer clmInputVector1 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector1); cl::Buffer clmInputVector2 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector2); cl::Buffer clmOutputVector = cl::Buffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pOutputVector);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

バッファーを作成するとき、コンテキスト（特定のデバイスではなく）、そのボリューム、および必要に応じてCL_MEM_COPY_HOST_PTRフラグを使用して、作成時にその中にコピーされるデータへのポインターが示されます。 前述したように、C ++ラッパーは参照カウントを使用するため、純粋なC APIとは異なり、手動でバッファーを削除する必要はありません。



次に、コードがファイル「OpenCLFile1.cl」に保存されているカーネルを作成する必要があります。 これを行うには、ファイルからテキストを読み取り、OpenCLプログラムを作成してコンパイルし、「TestKernel」という名前のカーネルを取得します。これについては次のパートで説明します。

 cl::Program::Sources source(1, std::make_pair(sourceCode.c_str(), sourceCode.length()+1)); cl::Program program = cl::Program(context, source); program.build(contextDevices); cl::Kernel kernel(program, "TestKernel");
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

コンパイル時には、実行する予定のデバイスを指定する必要があります。この場合は、テスト用に選択された1つのデバイスですが、一度にすべて指定できます。 コンパイルフラグを渡すこともできますが、この例では渡しません。



次に、カーネルに渡す引数を設定する必要があります。 CUDAとは異なり、引数ごとに特別な関数（C ++ラッパーの場合はメソッド）を呼び出し、必要に応じて引数のサイズを指定する必要があります。

 int iArg = 0; kernel.setArg(iArg++, clmInputVector1); kernel.setArg(iArg++, clmInputVector2); kernel.setArg(iArg++, clmOutputVector); kernel.setArg(iArg++, DATA_SIZE);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

ここで、最も重要なこと、つまりカーネルの起動に進みます。

 queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(DATA_SIZE), cl::NDRange(128));
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

実際には、queue.enqueueNDRangeKernelはカーネル開始コマンドをコマンドキューに追加し、処理する要素の数とグループのサイズを設定します。 （別の記事で）グループについて個別に説明しますが、ここでは、すべての要素が常にグループに分割され、パフォーマンスがグループのサイズに大きく依存するという事実のみを説明します。 この場合、要素の数はDATA_SIZEで、グループサイズは128です。カーネルの実行中に、DATA_SIZE回（不明なシーケンスで、場合によっては同時に）起動され、起動されるたびに、どの要素が処理されるかに関する情報が送信されます。

enqueueNDRangeKernelはブロックしないため、カーネルの起動後、カーネルが終了するまで待機する必要があります。

 queue.finish();
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

実際、finishは2つのタスクを実行します。

1）すべてのコマンドをデバイスに送信します（enqueueNDRangeKernelを実行すると、ドライバーがコマンドを受信して​​キューに入れることを保証しますが、デバイス上での実行を保証するものではなく、カーネルが実際に起動するまでにかなり長い時間がかかることがあります）。

2）キュー内のすべてのチームの完了を待ちます。

最初の部分のみを実行したい場合、ブロックされていないプッシュ（clFlush）コマンドがありますが、ドライバーは強制的にキューからコマンドの実行を開始します。



計算を実行した後、費やした時間を計算し、次のコマンドを使用して計算結果をホストにアップロードします。

 queue.enqueueReadBuffer(clmOutputVector, CL_TRUE, 0, DATA_SIZE * sizeof(float), pOutputVector);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

2番目の引数に応じて、enqueueReadBufferはブロックする場合としない場合があります。 私たちの場合、ブロックしているため、finishを個別に呼び出す必要はありません。 構文は単純です。最初の引数は読む場所、4番目の引数は読む量、最後の引数は読む場所です。 また、入力バッファの先頭からのオフセットを設定するパラメータもあります。これは、ホスト上のOpenCLバッファにアドレス演算を使用できないため、最初にデータを読み取る必要がない場合に使用する必要があります。

パート4-OpenCLカーネルコード

そしてここで、OpenCLでコードの記述を開始する必要があります（コードを呼び出すのは難しいので...甘やかします:)）。 OpenCLFile1.clは次のようになります。

 #include "MathCode.cpp" __kernel void TestKernel( __global const float* pInputVector1, __global const float* pInputVector2, __global float* pOutputVectorHost, int elementsNumber) { //Get index into global data array int iJob = get_global_id(0); //Check boundary conditions if (iJob >= elementsNumber) return; //Perform calculations pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]); }
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

順番に：

まず、コードにMathCode.cppファイルを含めます。MathCode.cppファイルには、数学関数が含まれています。これは、以前に注意を払うことを求めたもので、ホストでの従来の計算に使用されるものと同じです。 ご覧のとおり、コードをコピーすることさえせず、数学コードで同じファイルを使用します。

次に、__ kernelキーワードでマークしたカーネルを作成します。 一部のカーネル引数には、__ globalキーワードも付いています。これは、ホストコードで作成したデバイスのグローバルメモリ内のバッファーであることを示しています。

カーネルコードで、処理する必要がある要素の数を取得します。

 int iJob = get_global_id(0);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

get_global_idパラメーターは次元を示します。処理される要素は1、2、または3次元の配列になる可能性があるためです。

次に、境界条件を確認します。

 if (iJob >= elementsNumber) return;
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

これは、処理する要素の数が常にグループのサイズの倍数である必要があるため、処理する数を超える場合があるためです。

そしてチェックした後、主な部分を実行します：計算、ホストとまったく同じ方法で：

 pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

パート5-パフォーマンスのテストと測定

それでは、アプリケーションを起動し、パフォーマンスを評価して、いくつかの結論を導き出しましょう。



2台のマシンでテストを実行したところ、興味深い結果が得られました。

ノートブック（CPU： Intel®Core™i7-820QM 、GPU： NVidia Quadro FX 2800M ）：

 Host: 959.256 ms CPU: 82.4163 ms (13.106X faster than host) GPU: 9.90836 ms (109.014X faster than host)
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

デスクトップ（CPU： Intel®Core™i7-2600 、GPU： NVidia GeForce GTX 580 ）：

 Host: 699.031 ms CPU: 27.7833 ms (25.159X faster than host) GPU: 2.06257 ms (338.897X faster than host)
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

それでは、結果の分析を始めましょう。結果は非常に印象的です。 ラップトップのGPUはホストよりも約110倍高速であり、デスクトップでは約340倍高速ですが、印象的な結果です。 あなたが私にスリッパを投げ始めて、そのような比較は正しくないと言って始める前に、私は本当にいくつかのトリックがありますが、それ以上ではないと言います。



まず、ここでは、デバイスとの間でデータをコピーする時間を考慮していません。 一方で、これは間違っています。すべてをコピーすることを考慮すると、それほどうれしそうに見えないからです。 一方、コピーは計算と同時に実行できます。データが既にデバイスにある場合は、コピーをまったく実行する必要はありません。 一般に、すべては明確なものではなく、特定のタスクに依存します。



次に、数学コードがどのように見えたか覚えていますか？ それを見なかった人にとっては、同じデータに対する多くの数学的な操作であり、単純なコピーアンドペーストと係数の数値の置き換えによって判明しましたが、最初はより簡単で、1行だけでしたテストを開始しましたが、結果はそれほど喜ばしくありませんでした。GPUは4〜5倍高速でした。 どうしてそう思いますか？ （修辞的な質問、あなたは考えることはできません:)）。 そして、すべてが単純であり、メモリのパフォーマンスに出くわしました。 後でそれを手に入れて、メモリとプロセッサのパフォーマンスの関係に関する記事を書くことを願っていますが、これは別の話です。この記事では、このカーネルで演算プロセッサのパフォーマンスのクリーンなテストを得たという事実にのみ興味があります。



これらの2つの点を考えると、GPUはCPU上の非並列コードよりも純粋な算術で実際に数百倍高速であると言えます。これは一般に、理論的なパフォーマンスの違いに対応します。 （もう一つの希望は、実数と別の記事の理論への対応を測定するために手を伸ばすことです）。



しかし、GPUはすぐにそれを考慮し、テストの結果、CPUがOpenCLコードをかなり速く実行することがわかりました。正確には、デフォルト設定でMSVC10によってコンパイルされた通常のコードよりも13倍および25倍高速です。 それがどうなって、これらの数字がどこから来たのかを理解しましょう。



両方のプロセッサには4つの実コアと8つの仮想コアが含まれており、OpenCLはすべてのコアを使用するように作られていますが、4Xよりもはるかに改善されています。 そしてここで、OpenCLの実装で自動ベクトル化のサポートを追加したIntelに感謝する必要があります。 コードを変更することなく、OpenCLは使用可能なものに応じてSSEまたはAVXを使用します。 128ビットのSSEがあり、AVXが256ビットで動作することを考えると、パフォーマンスはそれぞれ16倍と32倍に上がるはずです。 これは真実に近いですが、まだ完全に一致しているわけではありません。 そして、TurboBoostのような楽しいことを覚えておく必要があります。 これらのプロセッサは1.73 GHz / 3.06 GHz（ラップトップ）と3.4 GHz / 3.8 GHz（デスクトップ）の周波数で動作しますが、実際にはラップトッププロセッサの周波数は1.73から2.8に連続的にジャンプし、非常に強力に加熱されるため（ここでは、湾曲した冷却システムのためにデルに大きなカメオを投げる必要があります）、3.06GHzの周波数テストでは、実際に重要な時間はありません。 さらに、実際の結果は常に理論的に可能なものよりも少ないことを忘れてはなりません（原則としてデスクトップはより速く動作するはずです）が、見てわかるように、同じハードウェアでほぼ無料で25倍のパフォーマンス向上が得られます。

おわりに

この記事の目標は、OpenCLを使用した作業の詳細をすべて説明することではなく、すべてがそれほど複雑ではないこと（ここでは既にそれほど単純ではないことを既に書いています）を示すことであり、理想的な条件では非常に印象的なパフォーマンスを得ることができます。同じハードウェア上で、さらに、すべてのデバイスに同じコードを使用できます。ただし、これらはほぼ理想的な条件であり、常にそうとは限らないことを忘れないでください。



PS：コードをいじって、別のハードウェアでテストを確認したい人のために、プロジェクト（および組み立てられた実行可能ファイル）がgithubにあります。開始するには、ハードウェアのメーカーのOpenCL SDKが必要になる場合があります。



PS2：誰もがIvy Bridgeを持っている場合、統合されたビデオコアのテストを見るのは興味深いでしょう。実際、OpenCL SDKの最新バージョンでは、IntelはIGPへのアクセスをオープンしましたが、最新世代のプロセッサのみを対象にしていますが、手元にはありません。AMDの結果は興味深いものです。