5.4.5.7. depth属性を使用するバッファーされたチャネルの実装

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

ダウンロード PDF

ID 683846

日付 4/01/2019

バージョン 19.1

Public

5.4.5.7. depth属性を使用するバッファーされたチャネルの実装

カーネルプログラムは、バッファーされたチャネルとバッファーされていないチャネルを含むことが可能です。チャネルの読み取り動作と書き込み動作が不均衡な場合、チャネル宣言にdepth属性を含めバッファーされたチャネルを作成し、カーネルがストールするのを防ぎます。バッファーされたチャネルは、異なるカーネルで並行して実行されているワークアイテムの動作を切り離します。

バッファーされたチャネルは、スループットの制限や、共有メモリーへのアクセスの同期化といった、データ・トラフィックの制御に使用することが可能です。バッファーされていないチャネルでは、読み取り動作がデータの値を読み取るまで書き込み動作を開始できません。バッファーされたチャネルでは、データの値がバッファーにコピーされるまで書き込み動作を開始できません。バッファーがフルの場合は、読み取り動作がデータの一部を読み取り、それをチャネルから削除するまで動作を開始することができません。

チャネルに対する消費率と生産率の一時的な不一致が予想される場合、depthチャネル属性を使用し、バッファーサイズを設定します。

次の例は、インテル® FPGA SDK for OpenCL™ のチャネル拡張を実装するカーネルコードでの、depthチャネル属性の使用方法を示しています。depth(N) 属性は、バッファーされたチャネルの最小深度を指定します。このNは、データ値の数を表します。

channel int c __attribute__((depth(10)));

__kernel void producer (__global int * in_data)
{
    for (int i = 0; i < N; i++)
    {
        if (in_data[i])
        {
            write_channel_intel(c, in_data[i]);
        }
    }
}

__kernel void consumer (__global int * restrict check_data,
                        __global int * restrict out_data)
{
    int last_val = 0;

    for (int i = 0; i < N, i++)
    {
        if (check_data[i])
        {
            last_val = read_channel_intel(c);
        }
        out_data[i] = last_val;
    }
}

この例において書き込み動作は、ブロッキングをすることなく10個のデータ値をチャネルに書き込むことができます。チャネルがフルになると、関連する読み取り動作がチャネルに発生するまで、書き込み動作を進めることはできません。

チャネルの読み取りと書き込みの呼び出しは条件付きステートメントのため、チャネルの読み取りと書き込みの呼び出しが不均衡になる可能性があります。チャネルにバッファー容量を追加することで、producerカーネルとconsumerカーネルを確実に分離することができます。この手順は、consumerカーネルがチャネルからデータを読み取っていないときに、producerカーネルがデータを書き込んでいる場合に特に重要です。

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

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

5.4.5.7. depth属性を使用するバッファーされたチャネルの実装