5.4.5.6. インテル® FPGA SDK for OpenCL™ チャネル実装を活用したモデル例

インテル® FPGA SDK for OpenCL™プロ・エディション: プログラミング・ガイド

ダウンロード PDF

ID 683846

日付 4/01/2019

バージョン 19.1

Public

5.4.5.6. インテル® FPGA SDK for OpenCL™ チャネル実装を活用したモデル例

カーネルの同時実行は、FPGAのチャネルの効率を向上させることができます。

同時実行を実現するため、ホストはカーネルを並列して起動します。該当する場合、カーネルはチャネルを介して互いに通信できます。

次に示すモデルは、安全かつ効率的な同時実行の活用方法の概要を示しています。

フィードフォワードのデザインモデル

フィードフォワードのデザイン・モデルを実装し、サイクルを作成せずに1つのカーネルから次のカーネルにデータを送信します。次のコードを参照ください。

__kernel void producer (__global const uint * src,  const uint iterations)
{
 for (int i = 0; i < iterations; i++) 
 {
  write_channel_intel(c0, src[2*i]);
  
  write_channel_intel(c1, src[2*i+1]);
 }
}

__kernel void consumer (__global uint * dst,  const uint iterations)
{
 for (int i = 0; i < iterations; i++) 
 {
  dst[2*i] = read_channel_intel(c0);
  
  dst[2*i+1] = read_channel_intel(c1);
 }
}

producerカーネルは、データをチャネルc0とc1に書き込みます。consumerカーネルは、c0とc1からデータを読み取ります。下の図は、2つのカーネル間のフィードフォワードのデータフローを表しています。

図 8. フィードフォワードのデータフロー

バッファー管理

フィードフォワードのデザインモデルでは、データはproducerとconsumerのカーネル間を一度に1ワードずつ移動します。複数のワードから構成される大きなデータメッセージの転送を容易にするため、通信用アプリケーションで一般的に見られるデザインパターンである、ピンポンバッファーを実装することができます。次の図は、カーネルとピンポンバッファー間の通信を表しています。

図 9. バッファー管理をともなうフィードフォワードのデザインモデル

managerカーネルは、producerカーネルとconsumerカーネル間における循環バッファーの割り当ておよび割り当て解除を管理します。consumerカーネルがデータを処理した後、managerはconsumerが解放したメモリー領域を受け取り、再度使用するためにproducerに送信します。managerはまた、使用されていない位置の初期セット (トークンの初期セット) をproducerカーネルに送信し、そこにproducerがデータを書き込めるようにします。

次の図は、バッファー管理の際に発生するイベントのシーケンスを表しています。

図 10. バッファー管理時のカーネルの通信

Managerカーネルは、トークンのセットをproducerカーネルに送信し、メモリー内のどの領域が現在使用されておらず、producerが利用できるかを示します。
managerがメモリー領域を割り当てた後、producerはピンポンバッファーのその領域にデータを書き込みます。
producerは書き込み動作完了後、consumerカーネルに同期トークンを送信し、処理するデータが含まれるメモリー領域を示します。次にconsumerカーネルは、ピンポンバッファーの該当領域からデータを読み取ります。
注: producer、consumer、managerカーネルは並行して実行されるため、consumerが読み取り動作を実行している間、producerは他の使用されていないメモリー位置へ処理を行うデータを書き込むことができます。
consumerは読み出し動作が完了後、メモリー領域を解放しトークンをmanagerに送り返します。次にmanagerカーネルはその領域をリサイクルし、producerが使用できるようにします。

OpenCLカーネルへのバッファー管理の実装

SDKが適切なバッファー管理を実行するためには、チャネルの読み取りおよび書き込みの順序が重要です。次のカーネル例を参照ください。

__kernel void producer (__global const uint * restrict src,
                        __global volatile uint * restrict shared_mem, 
                        const uint iterations)
{
    int base_offset;
    
    for (uint gID = 0; gID < iterations; gID++) 
    {
        // Assume each block of memory is 256 words
        uint lID = 0x0ff & gID;

        if (lID == 0)
        {
            base_offset = read_channel_intel(req);
        }

        shared_mem[base_offset + lID] = src[gID];

        // Make sure all memory operations are committed before
        // sending token to the consumer
        mem_fence(CLK_GLOBAL_MEM_FENCE | CLK_CHANNEL_MEM_FENCE);

        if (lID == 255)
        {
            write_channel_intel(c, base_offset);
        }
    }
}

このカーネルにおいて以下のコード行は独立しているため、インテル® FPGA SDK for OpenCL™オフライン・コンパイラーは、これらを同時に実行するようスケジュールすることができます。

shared_mem[base_offset + lID] = src[gID];

および

write_channel_intel(c, base_offset);

base_offsetにデータを書き込み、base_offsetをチャネルへ書き込むことは、グローバルメモリーへデータを書き込むよりもはるかに早い可能性があります。consumerカーネルは次に、チャネルからbase_offsetを読み取り、それをグローバルメモリーから読み取るためのインデックスとして使用します。同期がなければ、shared_mem[base_offset + lID] = src[gID];の実行が終了する前に、consumerがproducerからデータを読み取る可能性があります。その結果、consumerは無効なデータを読み取ることになります。このシナリオを回避するには、producerカーネルがデータをメモリーに収納した後に、同期トークンが発生しなければなりません。つまりconsumerカーネルは、producerがグローバルメモリーにデータを正常に格納するまでproducerカーネルからデータを消費することはできません。

この順序を維持するには、OpenCL mem_fenceトークンをカーネルに含めます。mem_fence構造は、CLK_GLOBAL_MEM_FENCEとCLK_CHANNEL_MEM_FENCEの2つのフラグを持ちます。mem_fenceは、mem_fence呼び出し前後に発生する動作間に、制御フローの依存性を効率的に作成します。CLK_GLOBAL_MEM_FENCEフラグは、グローバルメモリー動作が制御フローに従う必要があることを示します。CLK_CHANNEL_MEM_FENCEは、チャネル動作が制御フローに従う必要があることを示します。そのため、この例にあるwrite_channel_intel呼び出しは、グローバルメモリー動作が共有メモリーバッファーに格納されるまで開始できません。

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

インテル® FPGA SDK for OpenCL™プロ・エディション: プログラミング・ガイド

5.4.5.6. インテル® FPGA SDK for OpenCL™ チャネル実装を活用したモデル例

フィードフォワードのデザインモデル

バッファー管理

OpenCLカーネルへのバッファー管理の実装