5.10.1. シフトレジスターの推論

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

ダウンロード PDF

ID 683846

日付 4/01/2019

バージョン 19.1

Public

5.10.1. シフトレジスターの推論

シフトレジスターのデザインパターンは、多くのアプリケーションを効率的にFPGAへ実装するために非常に重要なデザインパターンです。しかし、シフトレジスターのデザインパターンを実装することに対し最初は違和感を抱くかもしれません。

次のコードを例に示します。

channel int in, out;

#define SIZE 512
//Shift register size must be statically determinable

__kernel void foo()
{
    int shift_reg[SIZE];
	   //The key is that the array size is a compile time constant

    // Initialization loop
    #pragma unroll
	for (int i=0; i < SIZE; i++)
	{
        //All elements of the array should be initialized to the same value
	    shift_reg[i] = 0;
    }
	
	while(1)
    {
        // Fully unrolling the shifting loop produces constant accesses
        #pragma unroll
        for (int j=0; j < SIZE–1; j++)
	    {
            shift_reg[j] = shift_reg[j + 1];
        } 
        shift_reg[SIZE – 1] = read_channel_intel(in);

        // Using fixed access points of the shift register
        int res = (shift_reg[0] + shift_reg[1]) / 2;

        // ‘out’ channel will have running average of the input channel
        write_channel_intel(out, res);
	}
}

各クロックサイクルで、カーネルは新しい値を配列にシフトします。このシフトレジスターをブロックRAMに配置することにより、インテル® FPGA SDK for OpenCL™オフライン・コンパイラーは、配列への複数のアクセスポイントを効率的に処理できます。シフトレジスターのデザインパターンは、フィルター (例えば、Sobelフィルターなどのイメージフィルターや、有限インパルス応答 (FIR) フィルターなどの時間遅延フィルター) を実装する際に理想的な方法です。

カーネルコードにシフトレジスターを実装する際は、次の点に注意してください。

シフトループを展開し、配列のすべての要素にアクセスできるようにします。
すべてのアクセスポイントは、一定のデータアクセスを持つ必要があります。例えば、複数のアクセスポイントを使用しネスト化されたループに計算を書き込む場合は、これらのループを展開し、一定のアクセスポイントを確立します。
配列の要素すべてを同じ値に初期化します。特定の初期値が不要な場合は、要素を初期化せずに維持することも可能です。
大規模な配列へのアクセスが静的に推論できない場合、オフライン・コンパイラーは非効率なハードウェアを作成することになります。それらのアクセスが必要な場合は、__privateメモリーの代わりに__localメモリーを使用してください。
大規模なシフトレジスターを条件付きでシフトしないでください。非効率なハードウェアの作成を防ぐため、シフトはシフトコードを含むループの反復で必ず行わなければなりません。

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

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

5.10.1. シフトレジスターの推論