5.5.5.5. depth属性を使用するバッファーされたパイプの実装

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

ダウンロード PDF

ID 683846

日付 4/01/2019

バージョン 19.1

Public

5.5.5.5. depth属性を使用するバッファーされたパイプの実装

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

バッファーされたパイプを使用し、スループットの制限や共有メモリーのアクセスの同期化といったデータのトラフィックを制御できます。バッファーされていないパイプにおいて書き込み動作は、読み出し動作がデータの読み取りをしようとしている場合にのみ開始できます。バッファーされていないパイプは、並行して実行されるカーネルでのブロッキングの読み書き動作と組み合わせて使用してください。バッファーされていないパイプは、自己同期型のデータ転送を効率的に提供します。

バッファーされたパイプにおいて書き込み動作は、受信するパケットを収容する容量がパイプにある場合にのみ進めることが可能です。読み出し動作は、少なくとも1つのパケットがパイプになければ実行することができません。

パイプ呼び出しが書き込みカーネルと読み出しカーネルで異なって表されている場合にバッファーされたパイプを使用すると、カーネルは並行して実行されません。

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

次の例は、OpenCLパイプを実装するカーネルコードにおけるdepth属性の使用方法を示しています。depth(N)属性は、バッファーされたパイプの最小深度を指定します。このNは、データ値の数です。読み出しカーネルと書き込みカーネルが、特定のバッファーされたパイプに異なる深度を指定する場合、インテル® FPGA SDK for OpenCL™オフライン・コンパイラーは2つの深度のうち大きい深度を使用します。

__kernel void
producer (__global int *in_data,
          write_only pipe int __attribute__((blocking))
		                    __attribute__((depth(10))) c)
{ 
    for (i = 0; i < N; i++)
    {
        if (in_data[i])
        {
            write_pipe( c, &in_data[i] );
        }
    }
}

__kernel void
consumer (__global int *check_data,
          __global int *out_data,
          read_only pipe int __attribute__((blocking)) c ) 
{
    int last_val = 0;
    for (i = 0; i < N; i++)
    {
        if (check_data[i])
        {
            read_pipe( c, &last_val );
        }
        out_data[i] = last_val;
    }
}

この例において書き込み動作は、10個のデータ値をパイプに正常に書き込むことができます。パイプがフルの状態になると書き込みカーネルは、読み出しカーネルがパイプのデータの一部を消費するまで失敗を返します。

パイプの読み出しと書き込みの呼び出しは条件付きステートメントのため、パイプの読み出しおよび書き込みの呼び出しには不均衡が発生する可能性があります。パイプにバッファー容量を追加すると、producerとconsumerカーネルを切り離すことができます。この方法は、consumerカーネルがパイプからデータを読み取っていない際に、producerカーネルがパイプにデータを書き込んでいる場合に特に重要です。

言語の選択

Intel.com 検索の使用

クイックリンク

最近の検索

詳細検索

検索のみ

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

5.5.5.5. depth属性を使用するバッファーされたパイプの実装