5.5.5.7. パイプ呼び出し順序の強制

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

ダウンロード PDF

ID 683846

日付 4/01/2019

バージョン 19.1

Public

5.5.5.7. パイプ呼び出し順序の強制

パイプの呼び出し順序を強制するには、メモリーフェンスもしくはバリアー機能をカーネルプログラムに導入し、メモリーアクセスを制御します。メモリーフェンス機能は、フェンス前後におけるパイプの同期呼び出し間に、制御フローの依存性を作成するために必要です。

インテル® FPGA SDK for OpenCL™オフライン・コンパイラーが計算ユニットを生成する際、それぞれが独立している命令のすべてに対し、命令レベルの並列性を構築するわけではありません。そのため、パイプの読み出しおよび書き込み動作は、それらに制御やデータの依存性がない場合でも、互いに独立して実行されない可能性があります。パイプ呼び出しが相互に通信する場合や、パイプが外部デバイスにデータを書き込む場合に、デッドロックが発生する可能性があります。

次のコード例は、producerカーネルとconsumerカーネルで構成されています。パイプc0とc1はバッファーされていないパイプです。c0とc1からのパイプ読み出し動作のスケジュールは、c0とc1へのパイプ書き込み動作と逆の順序になる可能性があります。つまり、producerカーネルが最初にc0に書き込む一方で、consumerカーネルはc1から読み出す可能性があるということです。consumerカーネルが空のパイプから読み出しているため、このパイプ呼び出しにおけるスケジューリングの変更は、デッドロックを発生させる可能性があります。

__kernel void producer (__global const uint * restrict src,
                        const uint iterations,
                        write_only pipe uint __attribute__((blocking)) c0,
                        write_only pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++) {
        write_pipe (c0, &src[2*i  ]);
        write_pipe (c1, &src[2*i+1]); 
    }
}

__kernel void consumer (__global uint * restrict dst,
                        const uint iterations,
                        read_only pipe uint __attribute__((blocking)) c0,
                        read_only pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++) {
        read_pipe (c0, &dst[2*i+1]);
        read_pipe( c1, &dst[2*i]); 
    }
}

パイプ呼び出しの順序を強制しデッドロックが発生しないようにするには、メモリーフェンス関数 (mem_fence) をカーネルに含めます。

各カーネルのパイプフラグとともにmem_fence呼び出しを挿入すると、書き込みおよび読み出しのパイプ呼び出しにシーケンシャルな順序付けが適用されます。次に、変更後のproducerとconsumerカーネルコードを示します。

__kernel void producer (__global const uint * src,
                        const uint iterations,
                        write_only_pipe uint __attribute__((blocking)) c0,
                        write_only_pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++)
    {
        write_pipe(c0, &src[2*i  ]);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        write_pipe(c1, &src[2*i+1]);
    }
}

__kernel void consumer (__global uint * dst;
                        const uint iterations,
                        read_only_pipe uint __attribute__((blocking)) c0,
		              read_only_pipe uint __attribute__((blocking)) c1)
{
    for(int i = 0; i < iterations; i++)
    {
        read_pipe(c0, &dst[2*i  ]);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        read_pipe(c1, &dst[2*i+1]);
    }
}

この例では、producerカーネルのmem_fenceは、c0へのパイプ書き込み動作がc1より先に発生するようにしています。同様にconsumerカーネルのmem_fenceは、c0からの読み取り動作がc1より先に行われるようにしています

パイプ使用時におけるカーネル間のメモリーの一貫性の定義

OpenCL™ Specification version 2.0 によると、カーネルの実行が完了しない限りメモリーの動作は定義されません。カーネルの実行は、メモリー動作に発生した変更が他のカーネルから見えるようになる前に終了させる必要があります。ただし、パイプを使用するカーネルは、共通のグローバル・メモリー・バッファーと、同期化されたメモリーアクセスを介しデータを共有できます。メモリーフェンスが渡された後にパイプに書き込まれたデータが読み出しパイプから確実に見えるようにするには、メモリーフェンスに関するメモリーの一貫性をカーネル間に定義します。

パイプの同期呼び出しとメモリー動作における制御フローの依存性を作成するには、mem_fence呼び出しにCLK_GLOBAL_MEM_FENCEフラグを追加します。

例:

__kernel void producer (__global const uint * restrict src,
                        const uint iterations,
                        write_only pipe uint __attribute__((blocking)) c0,
                        write_only pipe uint __attribute__((blocking)) c1)
{
    for (int i = 0; i < iterations; i++)
    {
        write_pipe(c0, &src[2*i]);
        mem_fence(CLK_CHANNEL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
        write_pipe(c1, &src[2*i+1]);
    }
}

このカーネルでmem_fence関数は、c0への書き込み動作とsrc[2*i] へのメモリーアクセスが、c1への書き込み動作とsrc[2*i+1] へのメモリーアクセスの前に必ず実行されるようにしています。これにより、c0に書き込まれたデータは、c1にデータが書き込まれる前に読み出しパイプから見えるようになります。

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

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

5.5.5.7. パイプ呼び出し順序の強制

パイプ使用時におけるカーネル間のメモリーの一貫性の定義