アルテラ+ OpenCL：コアを開く

みなさんこんにちは！



前の記事で、アルテラのFPGAで簡単なOpenCLの例を実行しました。

// ACL kernel for adding two input vectors __kernel void vector_add( __global const uint *restrict x, __global const uint *restrict y, __global uint *restrict z ) { // get index of the work item int index = get_global_id(0); // add the vector elements z[index] = x[index] + y[index]; }
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

私は意図的に詳細に立ち入らず、氷山の一角を示しました。開発プロセス、プロジェクトの組み立て、システムの起動です。



最初の記事を準備するとき、これらの行が何に変わるかについて非常に興味がありました（FPGA側から）。 アーキテクチャを理解することで、何かを最適化して、どのリソースが費やされているか、またこのシステムの良い点と悪い点を理解することができます。



この記事では、コアを開いて、次の質問に対する答えを見つけようとします。

• 彼のアーキテクチャは何ですか？
• どのように構成されていますか？ データはどのように処理されますか？
• どの周波数で機能しますか？ これはどのように決定されますか？
• RTLシミュレーターでコアのみをシミュレートすることは可能ですか？
• どのブロックが最もリソースを消費しますか？ これを最適化する方法はありますか？

その内部を見てみましょう！ 猫へようこそ！

アルテラの見方

FPGAのプロジェクトを詳しく検討する前に、ベンダーからのさまざまなプレゼンテーションに目を向けます。高レベル（マーケティング）言語での実装について話しているものです。



アルテラのOpenCLの大きなプレゼンテーション紹介をご覧になることをお勧めします

アルテラのOpenCL Compilerを使用したFPGAのパワーの活用 （ 注意 、100枚以上のスライド、約16 MB）。







ファームウェアの構成は次のとおりです。

• 周辺機器（PCIe、外部メモリ（DDR、QDR））へのアクセスを提供するIPコア。
• Nucleiはコンベアベースで実装されています。 OpenCLカーネルで説明されている計算を実行します。
• インフラストラクチャ： グローバルおよびローカルメモリインターコネクト 。

相互接続は、マスターとスレーブ（マスターとスレーブ）であるモジュール間の共通バスの分割です。



私たちの場合、ウィザードは、グローバルメモリ（ホストメモリまたは外部メモリのいずれか）と、キャッシュと呼ばれるローカル（内部）メモリの両方にデータを読み書きするカーネルです。 調停とデータ多重化のプロセスの結果、モジュールは、以下で説明するように、かなりの量のリソースを消費する可能性があります。



便宜上、モジュール間の通信プロトコルは標準化されています。 アルテラのプロジェクトでは、 Avalon ： Avalon-MM （Memory Mapped）やAvalon-ST （Streaming）などのインターフェースを使用しています。 これについては詳しく説明しません。読者はここで個別に読むことができます。 この記事では、相互接続のほとんどがAvalon-MMインターフェースになります。



繰り返しますが、これらはすべてOpenCLのカーネル記述から自動的に取得されることを強調します。

更新結果の投稿

前の記事で、Quartusバージョン14.1での作業に基づいたビルド結果について説明しました。

少し前、バージョン15.1がリリースされました。そこで大きな違いがあるかどうかを確認することにしました。 これを行うために、ソースコードを再生成し、新しいバージョンで再構築しました。



残念ながら、OpenCLのビジュアライザーとプロファイラーでは、（外観の）変更は発生していません。それらの外観には、まだ多くの要望があります。



--profileを使用したアセンブリレポート （プロファイリングカウンターを使用）：

 +-----------------------------------------------------------------------------------+ ; Fitter Summary ; +---------------------------------+-------------------------------------------------+ ; Fitter Status ; Successful - Sun Nov 22 13:18:14 2015 ; ; Quartus Prime Version ; 15.1.0 Build 185 10/21/2015 SJ Standard Edition ; ; Family ; Cyclone V ; ; Device ; 5CSEMA5F31C6 ; ; Timing Models ; Final ; ; Logic utilization (in ALMs) ; 5,472 / 32,070 ( 17 % ) ; ; Total registers ; 10409 ; ; Total pins ; 103 / 457 ( 23 % ) ; ; Total block memory bits ; 127,344 / 4,065,280 ( 3 % ) ; ; Total RAM Blocks ; 44 / 397 ( 11 % ) ; ; Total PLLs ; 2 / 6 ( 33 % ) ; ; Total DLLs ; 1 / 4 ( 25 % ) ; +---------------------------------+-------------------------------------------------+
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

以前のバージョンのコンパイラと比較して、プロジェクトは約100 ALMを失いました。



プロファイリングカウンターを使用しないビルドレポートは次のとおりです。

 +-----------------------------------------------------------------------------------+ ; Fitter Summary ; +---------------------------------+-------------------------------------------------+ ; Fitter Status ; Successful - Sun Nov 22 13:51:21 2015 ; ; Quartus Prime Version ; 15.1.0 Build 185 10/21/2015 SJ Standard Edition ; ; Family ; Cyclone V ; ; Device ; 5CSEMA5F31C6 ; ; Timing Models ; Final ; ; Logic utilization (in ALMs) ; 4,552 / 32,070 ( 14 % ) ; ; Total registers ; 7991 ; ; Total pins ; 103 / 457 ( 23 % ) ; ; Total block memory bits ; 127,344 / 4,065,280 ( 3 % ) ; ; Total RAM Blocks ; 44 / 397 ( 11 % ) ; ; Total PLLs ; 2 / 6 ( 33 % ) ; ; Total DLLs ; 1 / 4 ( 25 % ) ; +---------------------------------+-------------------------------------------------+
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

ご覧のとおり、約1000個のALMがプロファイリングカウンターとロジックで占められており、それらを「読み取り」ます。

将来的には、この特定のレポートを使用して、所要量を分析します。

プロジェクトを最初に見る

プロジェクトがgithubに投稿されていることを思い出させてください。



プロジェクトファイルは単純です。top.qpf （QPF-Quartus Project File）、最も重要なモジュールtop.vで、実際にはシステムモジュールのインスタンスとLEDに表示される単純なカウンターが含まれています。

システム（4535 ALM）

systemはQsysを使用して自動生成されたモジュールです 。 Qsysは、さまざまなIPブロックを接続し、相互接続、1つの周波数から別の周波数への切り替えなどに必要なモジュールのコードを自動的に生成できるGUIベースのツールです。



モジュール：

• vector_add_system（2141 ALM）は、 vector_addカーネルで記述した内容を実装するモジュールです。
• acl_iface（2343 ALM） -より便利なアクセスとカーネルとの対話を提供するインフラストラクチャ。

インターフェース：

• avs_vector_add_cra-カーネル管理用のAvalon-MM。
• avm_memgmem0_port_0_0_rw -DDRメモリにアクセスするためのAvalon-MM。 データ幅は256ビットです。

acl_iface（2343 ALM）

モジュール：

• pll（0 ALM） ： config_clk （50 MHz、外部ジェネレーターからの）からpll_outclk0 （100 MHz）クロックを受信するPLL 。
• acl_kernel_clk（1057 ALM） ：別のPLL：クロックを生成し、それがカーネルに供給されます。 興味深いニュアンスがあります。後で詳しく説明します。
• acl_kernel_interface（439 ALM） ：カーネルとプロセッサー間の「相互作用」を提供します（制御インターフェースと割り込みを介して）。
• clock_cross_kernel_mem1（82 ALM） ：異なる周波数（ CDC ）で動作するインターフェースの「調整」に従事しています。
• hps（0 ALM） ：これはHPS（ハードプロセッサシステム）インスタンスです。 FPGA tkのロジックは必要ありません。 これがハードウェアコアです。

インターフェース：

• f2h_sdram0 -DDRメモリにアクセスするためのAvalon-MMインターフェイス。 データ幅は256ビットで、動作周波数はpll_outclk0 （100 MHz）です。
• h2f_lw -AXIインターフェイス。 これにより、CPU（ARM）は、カーネルなどの制御/ステータスレジスタを使用してシステムを制御および構成できます。

これらのモジュールの合計容量を合計すると、合計は収束しません。 事実、Qsysはデフォルトでは相互接続タイプのモジュールを表示しません。 それらを表示するには、[ システム ]メニューの[ システムをQsysインターコネクトで表示 ]をクリックします。 その後、 mm_interconnect_ *という形式のモジュールがあり、568個と195個のALMを占有していることがわかります。

vector_add_system（2141 ALM）

このモジュールのアーキテクチャはGUIで表示できません。どのように機能するかを理解するために、Verilogに飛び込みます。



ダイアグラムの例は次のようになります。

• vector_add_system_interconnect_ *（443 ALM） -avm_memgmem0_port_0_0_rwインターフェイスの調停と多重化を実行する相互接続モジュール
• LSU_X（235） 、 LSU_Y（239） -ベクトルのカーネルメモリからデータを減算します（それぞれカーネル引数xおよびy ）。
• LSU_Z（424 ALM） -計算結果をグローバルメモリに書き込みます（引数z ）。
• acl_id_iterator（228 ALM） 、 acl_work_group_dispatcher（149 ALM） -カーネルによって実行されるタスクを発行します（どの要素を処理する必要があるかを示します）。
• acl_kernel_finish_detector（144 ALM） -カーネルがいつ動作を終了したかを決定します。

注：

LSUモジュールは1つのモジュール（ lsu_top ）のインスタンスであり、 lsu_local_bb0_ld_ 、 lsu_local_bb0_ld__u0およびlsu_local_bb0_st_addという名前です 。 便宜上、より「人間的な」名前を付けました。 LSUについては、以下で詳しく説明します。



カーネルの仕組み：

• 設定はCRAを通じて行われ、処理が開始されます。
• LSU_XおよびLSU_Yは 、データを読み取り、グローバルメモリに要求を行うための「コマンド」を受け取ります。
• 読み取りデータは、両方のLSUからのデータが準備できるまでメモリ（ FIFO ）にバッファリングされます。
• データが両方のFIFOにあるとすぐに、パイプラインに送信され、加算が実行されます。
• 結果はLSU_Zに分類され、そこで機会がグローバルメモリに書き込まれるのを待ちます。
• 必要な数の要素が処理され、すべての結果がメモリに書き込まれると（保留中のエントリはありません）、 kernel_finish_detectorがトリガーされ、 kernel_irq割り込みが設定されます。

3つのLSUは、グローバルメモリにアクセスするための1つのインターフェースを求めて互いに競合することに注意することが重要です。これらは、Avalon-MMインターフェースのマスターです。



図でadd_pipelineとして指定したパイプラインは、実際には別のモジュールには収まりません。vector_add_basic_block_0モジュールのvector_add.vファイルに配置されているだけです。



2つの32ビット数の加算を実行する行自体は、次のようになります。

 assign local_bb0_add = (rstag_3to3_bb0_ld__u0 + rstag_3to3_bb0_ld_);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

この行から作成される論理要素は、すべての有用な作業を行います。

それ以外はすべて、データをこのロジックに導くインフラストラクチャです。

LSU（ロードストアユニット）

このカーネルの最も興味深いモジュールはLSUです。 仕組みを見てみましょう。



実際、 lsu_topは、 READおよびSTYLEパラメーターに応じて選択される他のlsu_ *モジュールのラッパーです。



すべての種類のうち、2つのみがあります。

• LSU_READ_STREAMING - LSU_X 、 LSU_Y （READ = 1、STYLE = "STREAMING"）
• LSU_WRITE_STREAMING - LSU_Z （READ = 0、STYLE = "STREAMING"）

LSU_READ_STREAMING

モジュールのパラメーターに注意してください。

 BURSTCOUNT_WIDTH = 5; MEMORY_SIDE_MEM_LATENCY = 89;
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

BURSTCOUNT_WIDTHは、信号幅avm_burstcountを示します。Avalon -MMインターフェイスを介して要求された場合、 バーストトランザクション中に読み取る必要があるワード数があります。



信号幅が5の場合、最大バースト値は16です。これは、 仕様から明らかです。

 The value of the maximum burstcount parameter must be a power of 2. A burstcount interface of width n can encode a max burst of size 2^(n-1). For example, a 4-bit burstcount signal can support a maximum burst count of 8. The minimum burstcount is 1.
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

これは、最大で1つの要求が16個の256ビットワードで読み取られることを意味します。つまり、4096 Kビットまたは128個の32ビット数です（正確に32ビット整数を追加します）。



MEMORY_SIDE_MEM_LATENCYは、 lsu_burst_read_masterのFIFOワードの数に影響します 。 このFIFOは、グローバルメモリからの読み取りデータをバッファリングするために使用されます。



彼女の単語数はどのように決定されますか：

 localparam MAXBURSTCOUNT=2**(BURSTCOUNT_WIDTH-1); // Parameterize the FIFO depth based on the "drain" rate of the return FIFO // In the worst case you need memory latency + burstcount, but if the kernel // is slow to pull data out we can overlap the next burst with that. Also // since you can't backpressure responses, you need at least a full burst // of space. // Note the burst_read_master requires a fifo depth >= MAXBURSTCOUNT + 5. This // hardcoded 5 latency could result in half the bandwidth when burst and // latency is small, hence double it so we can double buffer. localparam _FIFO_DEPTH = MAXBURSTCOUNT + 10 + ((MEMORY_SIDE_MEM_LATENCY * WIDTH_BYTES + MWIDTH_BYTES - 1) / MWIDTH_BYTES); // This fifo doesn't affect the pipeline, round to power of 2 localparam FIFO_DEPTH = 2**$clog2(_FIFO_DEPTH);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

 _FIFO_DEPTH = 16 + 10 + ((89 * 4 + 32 - 1)/32) = 39    ,   : FIFO_DEPTH = 64
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

結論 ：

256ビットの64ワードのバッファ（キャッシュ）が割り当てられます。



実際、この事実を修正するために、ソースを選択する必要はありませんでした。アセンブリのレポートのRAM Summaryセクションを見るだけです。 計算は正しいことが判明し、レポートには7つのM10Kブロックが使用されることが示されています。 7つのブロックは、予想される256ビット* 64 = 16 Kbpsではなく、10,240ビット* 7 = 70 Kbpsです。



なぜこれが起こったのですか？

FPGAでは、内部メモリはさまざまな構成が可能な多数の小さなブロックです。

ここで、M10Kユニットを構成する方法を確認できます（つまり、Cyclone Vファミリのチップの基盤を形成します）。



メモリブロックの最大ワード長は40ビットです。256ビットでワードを作成する必要がある場合は、256/40 = 6.4-> 7ブロックが必要です。 メモリ内のワード数は64であるため、各ブロックは64x40として構成され、残りの75％のメモリは単純に使用されません。



バーストサイズとキャッシュサイズの影響は何ですか？

• バーストが多いほど、1回のリクエストで読み取ることができますが、残りのメモリリクエストはブロックされます（外部メモリと通信したいマスターが3人います）。
• キャッシュが大きいほど、新しいデータチャンクの読み取り中に処理用の「予約」にあるデータが多くなります。 大きなキャッシュのマイナスは、リソースの消費を除いてわかりません。 この場合、ワード数が256のキャッシュを作成でき、同じ数のM10Kブロックが使用されます。

LSU_WRITE_STREAMING

着信32ビットデータ（加算の結果）は、順番にFIFOに入れられます。 それぞれにMAXBURSTCOUNTがダイヤルされると（このモジュールではこのパラメーターも16）、メモリへの書き込みトランザクションが発生します。 これらの各FIFOのデータ幅は32です。このようなFIFOは8個（256/32）あります。



これらのfifoはどのくらいのデータ向けに設計されていますか



コードlsu_streaming_writeから計算を行います。 このモジュールパラメータ用。 MEMORY_SIDE_MEM_LATENCYは32です。

 localparam MAXBURSTCOUNT=2**(BURSTCOUNT_WIDTH-1); localparam __FIFO_DEPTH=2*MAXBURSTCOUNT + (MEMORY_SIDE_MEM_LATENCY * WIDTH + MWIDTH - 1) / MWIDTH; localparam _FIFO_DEPTH= ( __FIFO_DEPTH > MAXBURSTCOUNT+4 ) ? __FIFO_DEPTH : MAXBURSTCOUNT+5; // This fifo doesn't affect the pipeline, round to power of 2 localparam FIFO_DEPTH= 2**($clog2(_FIFO_DEPTH));
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

 MAXBURSTCOUNT = 2^4 = 16 __FIFO_DEPTH = 2 * 16 + ( 32 * 32 + 256 - 1)/256 = 36 + 5 = 41 _FIFO_DEPTH = 41      : FIFO_DEPTH = 64
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

レポートで確認します：64 * 32 = 2048ビット（1 M10K）。



FIFOは完全に独立しているため、各FIFOに1つのM10Kブロックが割り当てられ、 lsu_read_streamingの 7つのM10Kブロックに対して8つのM10Kブロックになります。



なぜ8個のFIFOを作成したのですか？ ほとんどの場合、これは簡単です（有効な単語の数を個別に保存する必要はありません）。

LSUパラメータはどのように計算されますか？

そのような数字がどこから来たのかを理解してみましょう：

これらの設定は、ボードを説明するファイル（altera / 15.1 / hld / board / de1soc / de1soc_sharedonly / board_spec.xml）から取得された疑いがあります。



グローバルメモリに関連付けられている行を見つけます。

  <!-- One DDR3-800 DIMM, 256-bit data --> <global_mem max_bandwidth="6400"> <interface name="acl_iface" port="kernel_mem0" type="slave" width="256" maxburst="16" latency="240" address="0x00000000" size="0x40000000"/> </global_mem>
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

これらのパラメーターの説明については、 Altera SDK for OpenCL：カスタムプラットフォームツールキットユーザーガイドの章board_spec.xmlファイルのXML要素、属性、およびパラメーターを参照してください 。



max_bandwidth- 現在の構成で結合されたすべてのグローバルメモリインターフェイスの最大帯域幅。 アルテラのオフライン・コンパイラは、max_bandwidthを使用して、アプリケーションとボードに適したアーキテクチャを選択します。 思い出のデータシートからこの帯域幅の値を計算します。







残念ながら、どのユニットとどのように考えられているかについての説明はありません。一方で、プロファイラーで6400 MB / s が書き込まれ、6400 MB / sの計算によると、まったく機能しません：400（MHz、DDRクロック周波数）* 32（ビット、幅DDRインターフェイスのデータ信号）* 2（2つの前面で動作）= 25600 Mb / s = 3200 MB / s または、両方向で検討する必要がありますか？



max_burst- スレーブインターフェイスの最大バーストサイズ。





この例では-16で、BURSTCOUNT_WIDTH = 5になります。 fpga2hps_sdramインターフェースはmax_burstcount = 128をサポートします。16-これはある種のマジックナンバーで、誰にでも適していますか？ :)



レイテンシ - メモリインターフェイスがリクエストに応答するまでの時間をナノ秒（ns）で指定する整数。 レイテンシは、カーネルがボードシステムにメモリ読み取り要求を発行してから、カーネルに戻るメモリデータへのラウンドトリップ時間です。 たとえば、クロッククロッシングブリッジを備えた200 MHzで動作するアルテラDDR3メモリコントローラのレイテンシは約240 nsです。







この場合も、遅延が240 nsであるとします。 明らかに、ディメンションMEMORY_SIDE_MEM_LATENCYはティックの数です（このコメントが示唆するのは、LSUとメモリ間のサイクルのレイテンシです）。



board_spec.xmlの値（ maxburst 、 latency ）およびカーネルの構造（追加される引数の数（ readers ））を変更して、いくつかの実験を行ってみましょう 。 両方のモジュール（ LSU_X （lsu_read_streaming）およびLSU_Z （lsu_write_streaming））のMEMORY_SIDE_MEM_LATENCYパラメーターの値を監視します。

 |--------------------------------------------------------| | maxburst | latency | readers | MEMORY_SIDE_MEM_LATENCY | | | | |-------------------------| | | | | LSU_X | LSU_Z | |--------------------------------------------------------| | 16 | 0 | 1 | 25 | 16 | | 16 | 100 | 1 | 45 | 16 | | 16 | 240 | 1 | 73 | 16 | |--------------------------------------------------------| | 16 | 0 | 2 | 41 | 32 | | 16 | 100 | 2 | 61 | 32 | | 16 | 240 | 2 | 89 | 32 | |--------------------------------------------------------| | 16 | 0 | 3 | 57 | 48 | | 16 | 100 | 3 | 77 | 48 | | 16 | 240 | 3 | 105 | 48 | |--------------------------------------------------------| | 32 | 0 | 1 | 41 | 32 | | 32 | 100 | 1 | 61 | 32 | | 32 | 240 | 1 | 89 | 32 | |--------------------------------------------------------| | 32 | 0 | 2 | 73 | 64 | | 32 | 100 | 2 | 93 | 64 | | 32 | 240 | 2 | 121 | 64 | |--------------------------------------------------------| | 32 | 0 | 3 | 105 | 96 | | 32 | 100 | 3 | 125 | 96 | | 32 | 240 | 3 | 153 | 96 | |--------------------------------------------------------|
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

依存関係は何ですか：

• レイテンシが増加し、 maxburstおよびリーダーが 修正されると、 LSU_X_MEMORY_SIDE_MEM_LATENCYはレイテンシ/ 5に等しい値だけ増加します。 ほとんどの場合、5は5 nsです（200 MHzのマジック周波数に送信されていますか？）。
• リーダーの数が増えると、 LSU_X_MEMORY_SIDE_MEM_LATENCYはバーストの値だけ増加します。
• LSU_Z_MEMORY_SIDE_MEM_LATENCYは、グローバルメモリにアクセスする要素の数と最大バースト値に線形に依存します。

数式を見ることができます：

• LSU_X_MEMORY_SIDE_MEM_LATENCY = 9 + リーダー * maxburst + latency / 5。 （9は何らかのマジックナンバーか、私が最後まで到達していないカーネルの別の特性です。おそらくこれはカーネルの一般的な遅延です）。
• LSU_Z_MEMORY_SIDE_MEM_LATENCY = maxburst * readers 。

注 ：

これらの式は特定のカーネル（実装）専用であり、別の式ではすべてが異なる場合があります。

コアはどの周波数で機能しますか

コアクロックはacl_kernel_clkモジュールによって生成されます。

これは、動的に再構成（出力周波数を変更）できるPLLに基づいています。



Qsysまたはsystem_acl_iface_acl_kernel_clk_kernel_pll.vでこのモジュールを開くと、このPLLが140 MHz（ kernel_clk ）と280 MHz（ kernel_clk2x ）の2つの信号を生成することがわかります 。 kernel_clk2xはどこでも使用されていないと言わなければなりません。



コアは常に（および任意の）140 MHzでのみ動作し、オーバークロックできないことを意味しますか？ もちろん違います。

140 MHzは、この特定のボードの設定です。



使用される論理要素とそれらの接続方法に応じて、回路が障害なく動作することが保証されるクロック周波数の値は異なる場合があります。 この問題については、パイプライン処理に関する記事で触れました。



コンパイラのタスクは、特定の周波数要件を満たすようにプリミティブ（論理要素、メモリブロックなど）を配置することです。 これは次のことを意味します。

• 彼は最大クロック速度を与える場所を見つけようとしません。
• しばらくの間、彼がチップ内の要素の配置を整理し、繁殖できないことを理解した場合、彼は最良の選択肢の1つを残します（検索中）。

140 MHzの代わりに、 Quartusは135 MHzの最大クロック速度を示しているとします。 これは次のことを意味します。

• コンパイラは、135 MHzが適用されると、計算が正しく行われ、何もハングしないなどを保証します（もちろん、コード自体にアルゴリズムエラーがない場合）。
• 140 MHzを適用すると、すべてが正常になる場合があります。 またはそうでないかもしれません。 それはチップに依存します-同じラベル付けのチップはわずかに異なる場合があるため、コンパイラは再保険され、最悪の場合に計算されます。

ほとんどの場合、FPGAプロジェクトを再構築した後、開発者はアセンブリレポートを見て、周波数がスキームに適合するかどうか疑問に思っています。前回の記事では、バイナリを取得してステッチしました。コンパイラがこれらの140 MHzに適合しない場合はどうなりますか？計算は間違っていますか？



この問題を開発者から隠すために、アルテラは非常に興味深いトリックを作成しました（おそらく、アルテラのOpenCL SDKで遊んでいるときに掘り下げたものの中で最も興味深いものです）。

• ビルドが完了すると、adjust_plls.tclスクリプトが呼び出されます。最大許容コア周波数（Fmax）を受け取り、pll_romモジュールのROMを初期化するために使用されるファイル（pll_rom.mifおよびpll_rom.hex）を生成します。
• FPGAがロードされると、設定された周波数（140 MHz）がロジックに適用されます。カーネルを起動する前に、ROMからデータが読み取られ、これらの係数を使用して、PLLが再構築されます（再構成インターフェイスを介して）。再構成が終了するとすぐに、目的の周波数がすでにカーネルに供給されます。

合計：

• コアでは、周波数を適用できるものが提供されます。ロジックの容量が大きすぎることが判明し、指定された数値を満たせなかった場合、計算は中断されず、遅くなります。
• 許容周波数が高い場合、PLLはこの値に調整されます（計算が加速します）。最大頻度を与える場所の検索はありません。それでも「分散」できるという感覚がある場合は、PLL周波数で手動でバーを上げる方が良いでしょう。

アセンブリを少し簡素化する

カーネルがどのように配置および構成されているかを学習する前に、（カーネルに）変更を加えるか、ハードウェアでデバッグする場合に役立つ小さな余談を行います。開発プロセスを



思い出させてください。FPGAファームウェアを含むvector_add.aocxファイルはvector_add.clから取得されます。問題は、Quartusプロジェクトに変更を加えても、* .aocxに該当しないことです。aocユーティリティが再起動すると、「デフォルトプロジェクト」がコピーされ、Verilog IPが再生成されます。したがって、変更は失われます。Aocユーティリティ







はバイナリですが、呼び出されたときにトレースできます：

 $ aoc device/vector_add.cl -o bin/vector_add.aocx --board de1soc_sharedonly --profile -v
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

このスクリプトはaoc.pl pearl で実行されますが、これはすでにすべての有用な作業を行っています。aoc

ユーティリティを使用せずに、このスクリプトを直接呼び出すことができます。

 $ /home/ish/altera/15.1/quartus/linux64/perl/bin/perl /home/ish/altera/15.1/hld/share/lib/perl/acl/aoc.pl device/vector_add.cl --board de1soc_sharedonly --profile -v
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

スクリプトがインタプリタ言語で書かれているのは良いことです。つまり、スクリプトの機能を理解し、変更を加えることができます。

スクリプトの最初に、キー（ヘルプでユーザーに非表示になっているものを含む）を介して設定されるさまざまな変数が説明されています。



そのため、そこに--quartusキーが見つかります。これにより、quartusのアセンブリと、* .aocxファイル内の必要なパーツのパッケージ化のみが行われます。プロジェクト（ソース）の再生成はありません。



また、利便性を高めるために、コンソールにビルドログを表示できます。これを行うには、mysystem_full関数呼び出しでstdoutとstderrを空行として指定します。

 $return_status = mysystem_full( {'time' => 1, 'time-label' => 'Quartus compilation', 'stdout' => '', 'stderr' => ''}, $synthesize_cmd);
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

これで、プロジェクトに簡単に変更を加え（最適化を行い、SignalTapを追加）、clang呼び出しとコード再生成を使用してカーネル全体ではなく、FPGAのプロジェクトのみを再構築することができます。



これを確認するために、インターフェイスにSignalTapを追加しました（また、カーネルをロードして計算を開始した後、デバッガーを使用して接続できるように15秒のスリープを追加しました）。

カーネルの制御方法

avs_vector_add_craインターフェースを使用してカーネルを構成します。データはレジスターのアドレスに書き込まれます。



残念ながら、オープンアクセスのレジスタマップとその設定方法が見つからなかったため、少し調査する必要があります。



すべてのレジスタはvector_add.vで説明されており、適切な名前が付けられています。



これらは64ビットです。[31：0]は下位32ビットを示し、[63:32]-最上位を示します。

 0x0 - status 0x1 - 0x4 - profile 0x5 - [31:0] - work_dim 0x5 - [63:32] - workgroup_size 0x6 - [31:0] - global_size[0] 0x6 - [63:32] - global_size[1] 0x7 - [31:0] - global_size[2] 0x7 - [63:32] - num_groups[0] 0x8 - [31:0] - num_groups[1] 0x8 - [63:32] - num_groups[2] 0x9 - [31:0] - local_size[0] 0x9 - [63:32] - local_size[1] 0xA - [31:0] - local_size[2] 0xA - [63:32] - global_offset[0] 0xB - [31:0] - global_offset[1] 0xB - [63:32] - global_offset[2] 0xC - [31:0] - kernel_arguments[31:0] - input_x[31:0] 0xC - [63:32] - kernel_arguments[63:32] - input_x[63:32] 0xD - [31:0] - kernel_arguments[95:64] - input_y[31:0] 0xD - [63:32] - kernel_arguments[127:96] - input_y[63:32] 0xE - [31:0] - kernel_arguments[159:128] - input_z[31:0] 0xE - [63:32] - kernel_arguments[191:160] - input_z[63:32]
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

名前に基づいて、何かをランダムに構成および実行しようとすることができますが、それを危険にさらすのではなく、そこに何がどの順序で書かれているかを見つけてください。



このインターフェイスですべてのトランザクションを記録します（SignalTapを使用）：

 ---------------------------------------------- | addr | write_data | byte_enable | ---------------------------------------------- | 0x5 | 0x00000000 0x00000001 | 0x0F | | 0x5 | 0x000F4240 0x00000000 | 0xF0 | ---------------------------------------------- | 0x6 | 0x00000000 0x000F4240 | 0x0F | | 0x6 | 0x00000001 0x00000000 | 0xF0 | ---------------------------------------------- | 0x7 | 0x00000000 0x00000001 | 0x0F | | 0x7 | 0x00000001 0x00000000 | 0xF0 | ---------------------------------------------- | 0x8 | 0x00000000 0x00000001 | 0x0F | | 0x8 | 0x00000001 0x00000000 | 0xF0 | ---------------------------------------------- | 0x9 | 0x00000000 0x000F4240 | 0x0F | | 0x9 | 0x00000001 0x00000000 | 0xF0 | ---------------------------------------------- | 0xA | 0x00000000 0x00000001 | 0x0F | | 0xA | 0x00000000 0x00000000 | 0xF0 | ---------------------------------------------- | 0xB | 0x00000000 0x00000000 | 0x0F | | 0xB | 0x00000000 0x00000000 | 0xF0 | ---------------------------------------------- | 0xC | 0x00000000 0x20100000 | 0x0F | | 0xC | 0x00000000 0x00000000 | 0xF0 | ---------------------------------------------- | 0xD | 0x00000000 0x20500000 | 0x0F | | 0xD | 0x00000000 0x00000000 | 0xF0 | ---------------------------------------------- | 0xE | 0x00000000 0x20900000 | 0x0F | | 0xE | 0x00000000 0x00000000 | 0xF0 | ---------------------------------------------- | 0x0 | 0x00000000 0x00000001 | 0x0F | ----------------------------------------------
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

注：

byte_enableは、レジスタのどのバイトを書き込むかを「選択」します。たとえば、最初のトランザクションでは、0x00000001を0x5レジスタの下位32ビットに書き込みました（上位32ビットは変更されませんでした）。



SignalTapでトランザクションを監視することは必ずしも便利ではない場合があります。ホストでは、環境変数を使用して追加のデバッグを有効にできます。これらは、アルテラStratix Vネットワークリファレンスプラットフォームポーティングガイドのトラブルシューティングの 章で確認できます。ACL_HAL_DEBUG変数が必要です。我々は2にその値を公開して、ホストアプリケーションを実行vector_add：

 root@socfpga:~/myvectoradduint# export ACL_HAL_DEBUG=2 root@socfpga:~/myvectoradduint# ./vector_add // <  > :: Launching kernel 0 on accelerator 0. :: Writing inv image [ 0] @ 0x28 := 1 :: Writing inv image [ 4] @ 0x2c := f4240 :: Writing inv image [ 8] @ 0x30 := f4240 :: Writing inv image [12] @ 0x34 := 1 :: Writing inv image [16] @ 0x38 := 1 :: Writing inv image [20] @ 0x3c := 1 :: Writing inv image [24] @ 0x40 := 1 :: Writing inv image [28] @ 0x44 := 1 :: Writing inv image [32] @ 0x48 := f4240 :: Writing inv image [36] @ 0x4c := 1 :: Writing inv image [40] @ 0x50 := 1 :: Writing inv image [44] @ 0x54 := 0 :: Writing inv image [48] @ 0x58 := 0 :: Writing inv image [52] @ 0x5c := 0 :: Writing inv image [56] @ 0x60 := 20100000 :: Writing inv image [60] @ 0x64 := 0 :: Writing inv image [64] @ 0x68 := 20500000 :: Writing inv image [68] @ 0x6c := 0 :: Writing inv image [72] @ 0x70 := 20900000 :: Writing inv image [76] @ 0x74 := 0 :: Accelerator 0 reporting status 2. :: Accelerator 0 is done.
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

0x28は、5番目の64ビットレジスタのバイトオフセットであることに気付くかもしれません。

アドレスとデータが一致することがわかりますが、このデバッグでは、ゼロレジスタにトランザクションに関する情報がありません（ACL_HAL_DEBUGが 5に設定されている場合でも）。



設定結果：

• work_dim - 0x1-以降 1次元のベクトルがあります。
• workgroup_size -0xF4240または1,000,000。
• global_size-最初の次元の場合は0xF4240、その他すべての次元の場合は0x1。
• num_groups-すべての次元で0x1。
• local_size-最初の次元では0xF4240、その他すべての次元では0x1。
• global_offset-すべての次元で0x0。
• input_x、input_y、input_z - 0x20100000、0x20500000、0x20900000、それぞれ。

カーネルを調整した後、ゼロレジスタの最下位ビットが動いて、計算を開始します。

コアシミュレーション

これで、カーネルの構成方法がわかりました。シミュレートしましょう。

カーネルには、チューニング用とデータ読み取り用の2つのインターフェイスしかありません（割り込みを設定するためのインターフェイス（1つの信号サイズ）がありますが、あまり興味がありません）。



カーネルをシミュレートするには、実際と同じようにすべてを行う必要があります。

• それを設定します（書き込みの順序がわかり、インターフェースは比較的単純です）。
• バッファが割り当てられ、データが保存されるグローバルメモリへのアクセスを提供します。

もちろん、Linuxやホストアプリケーションをシミュレートすることは望ましくないため、最初の近似として、次のスキームに制限することができます。

• cra_driver-カーネルを構成するためのドライバー。
• vector_add_system-シミュレートするカーネル（DUT）。
• avalon_mm_clock_crossing-コア周波数（140 MHz）からコントローラーからの読み取り周波数（100 MHz）へ、またはその逆へのデータ転送。
• prepare_data-シミュレーションを開始する前にXおよびYバッファーにデータを書き込む単純なタスク。
• avalon_mm_interconnect -2つのAvalon-MMインターフェイスの多重化と調停。
• ddr3_contoller、ddr3_model-アルテラのハードメモリコントローラとDDR3メモリのシミュレーションモデル。モデルとコントローラーの設定は、hpsモジュールで使用される設定と同じです。

カーネルを構成するためのドライバーは、SignalTapのおかげで明らかになったカーネル設定を使用した、次のタスクの単純な呼び出しです。

 task cra_write( input bit [3:0] _addr, bit [63:0] _data, bit [7:0] _byteenable ); $display("%m: _addr = 0x%x, _data = 0x%x, _byteenable = 0x%x", _addr, _data, _byteenable ); @( posedge clk ); cra_addr <= _addr; cra_wr_data <= _data; cra_byteenable <= _byteenable; cra_wr_en <= 1'b0; @( posedge clk ); cra_wr_en <= 1'b1; @( posedge clk ); cra_wr_en <= 1'b0; // dummy waiting repeat (10) @( posedge clk ); endtask
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

計算を行う必要があるデータを準備するために、事前にパックされたアドレスでメモリに書き込まれる同様のタスクを作成します。 （アドレスの割り当てを誰か（システム）に依頼する必要はありません）。デモンストレーションでは、タスクの条件に応じて必要に応じて2x1000000の数字を書く必要はありません-動作を確認するには数千個で十分です。データを書き込まない場合、x（不明な値）がメモリから読み取られます。これらのセルには何も書き込まれませんでした。



シミュレーションの結果を確認します（16進数形式の一時的な建物のすべての数字）（スクリーンショットは別のウィンドウで開くことをお勧めします）：開始信号が来ます





そして、いくつかのストロークの後、両方のLSUは同時にberstaのサイズにデータを読み出すための要求を露出さが0x10に等しい= 16興味深いことに、3つの要求は、最初にのみ受け入れられLSU_X、次いでからLSU_Y：それは信号を示しavm_waitrequest場合にのみ受け入れられるリード要求をwaitrequestはゼロに等しい。ご想像のとおり、リクエストが交互に行われない理由は、シェダラーとグローバルメモリへの相互接続の問題です。



なぜ正確に3つの要求があったのですか？

LSUのFIFOのワード数は64であり、4つの読み取り要求が予想されますが、このようなトリックがありました。

 parameter READTHRESHOLD = FIFODEPTH - MAXBURSTCOUNT - 4; assign too_many_reads_pending = (reads_pending + fifo_used) >= READTHRESHOLD; // make sure there are fewer reads posted than room in the FIFO
      
      

        
        
        
      

    
        
        
        
      
      

        
        
        
      

    
     

読み取り要求を3つ送信しました（3 x 16 = 48ワード）：16ワードのスペースはまだありますが、4ワードが完了するまで読み取りは停止します。 12の小さなバーストはありません。このモジュールはその方法を知りません（その意味はありません-ロジックの無駄です）。



しばらく後に、データを読み出すようになっ（信号を参照。LSU_X_avm_readdatavalidとすぐ露出信号）LSU_X_o_valid、その中の32ビットデータ知らせるLSU_X_o_readdataは、さらなる処理のために準備ができて、私たちはコーキング：LSU_X_i_stallのいずれかに設定されています。実際、データを追加するコンベヤーにはLSU_Yからのデータがありません。したがって、LSU_Y_avm_readdatavalidが到着するまで、すべてがシャットダウンします。そして、LSU_Y_o_validを設定します。



これに続く次のメジャーは、有効性についてLSU_Z_i_writedataに通知するLSU_Z_i_valid です。追加パイプラインは1つのメジャーで機能しました。LSU_Zは、必要な量のデータの蓄積を待機しています。これは、もちろん16 256ビットワードであり、書き込みトランザクションを作成します。同時に、LSU_XとLSU_Yは、FIFOが空になると読み取りトランザクションを実行します。LSU_X_i_stallとLSU_Y_i_stallはコックされていないことに注意してください。これは、パイプラインがシャットダウンせず、各クロックが新しいデータを受信することを意味します。













多くのトランザクションを調べた結果、より多くのギャップが発生しないことは明らかです。



結論：

パイプラインが停止することはない（最初の段階を除いて）ため、最大のパフォーマンスで動作することは明らかであり、この単純な例では狭いリンクが彼です。読み取りは256ビットワードを使用して行われるため、8つの32ビット数の並列加算を整理することは理にかなっています。その場合、ボットネクはメモリアクセスで発生する可能性が高くなります。使用に加算器のプロシージャ・EED番号特別な意味はありません



シミュレーションのおかげで、RTLコードにいくつかの変更（バースト長やキャッシュサイズなど）を加えて、これが計算の速度にどのように影響するかを確認し、非常に迅速に答えを確認できます。これを10〜15分待つ必要はありませんQuartusでプロジェクト全体を再構築するには、シミュレータでシミュレーションを実行します。



もちろん、* .clファイルを編集して（たとえば、特別なディレクティブを使用して）アーキテクチャを変更する必要があります。コードを再生成し、シミュレーションを実行して、プロジェクト全体を組み立てずに結果のゲインを確認することもできます。

おわりに

vector_addカーネルがFPGA側から何に変わるか、そしてどのように構成されているかを学びました。



ご覧のとおり、FPGA開発者が抱えていた多くの問題（ゼロから作成した場合）は解決されました。

• キャッシュサイズ
• 建築
• インターフェイスの相互接続と調停を行いました
• 結果の値へのクロック周波数の自動調整

もちろん、自動的に行われることが最適であるという事実ではありません。





低レベルの最適化と手動チューニングが最後のステップです。gccをビルドした後、アセンブラーに頻繁に入りますか？まず、* .clで高度な最適化を実行し、Quartusの設定で遊ぶ必要があります。



コアをシミュレートする機会があるのは良いことです。手元に鉄（ボード）がなくても、コアのパフォーマンスを推定できます。



ご清聴ありがとうございました！ コメントまたは個人のメールで質問やコメントをお待ちしております。