5.4.5.8. チャネルの呼び出し順序の強制

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

ダウンロード PDF

ID 683846

日付 4/01/2019

バージョン 19.1

Public

インテルのみ表示可能 — Ixiasoft

詳細を見る

5.4.5.8. チャネルの呼び出し順序の強制

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

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

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

__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++)
    {
        /*During compilation, the AOC might reorder the way the consumer kernel 
        writes to memory to optimize memory access. Therefore, c1 might be read
        before c0, which is the reverse of what appears in code.*/  

        dst[2*i+1] = read_channel_intel(c0);
        dst[2*i] = read_channel_intel(c1);
    }
}

デッドロックの発生を防ぐには、カーネルにメモリーフェンス関数 (mem_fence) を含め、チャネル呼び出しの順序を強制します。

各カーネルのチャネルフラグとともにmem_fence呼び出しを挿入すると、書き込みと読み取りのチャネル呼び出しに、シーケンシャルな順序付けが適用されます。次は、変更されたproducerとconsumerカーネルのコードです。

channel uint c0 __attribute__((depth(0)));
channel uint c1 __attribute__((depth(0)));

__kernel void producer (__global const uint * src,
                        const uint iterations)
{
    for (int i = 0; i < iterations; i++)
    {
        write_channel_intel(c0, src[2*i]);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        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+1] = read_channel_intel(c0);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        dst[2*i] = read_channel_intel(c1);
    }
}

この例でproducerカーネルのmem_fenceは、c0へのチャネル書き込み動作がc1よりも前に発生するようにしています。同様にconsumerカーネルのmem_fenceは、c0のチャネル読み取り動作がc1よりも前に発生するようにしています。

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

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

チャネルの同期呼び出しとメモリー動作間の制御フローの依存性を作成するには、CLK_GLOBAL_MEM_FENCEフラグをmem_fenceコールに追加します。

例:

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

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

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

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

5.4.5.8. チャネルの呼び出し順序の強制

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