スーパーコンピュータでも使われ始めた FPGA (4)

このコースでは、スーパーコンピュータ (以下、ブログの文字数節約のために一般的に使われている略称「スパコン」を使います) にも FPGA が何故使われ始めたのか、FPGA をスパコンに使うことによって何ができるようになるのか、どのような技術的な面白さがあるのか等について紹介していきます。

第4回のこの記事では、FPGA を導入した国内スパコンである Cygnus を対象にどのような研究開発が行われているかについて紹介していきます。

OpenCL から制御可能なデータ転送技術の開発

前回紹介した Cygnus は、複数の GPU と FPGA とを協調動作させてアプリケーションを加速させることを狙ったスパコンです。しかし、このコンセプトは言うは易く行うは難しで、この狙いを実現するためには解決しなければならない課題がいくつもあります。

今回の記事では、その課題の一つである「計算科学の研究者・開発者がなるべく利用しやすく、かつ高性能なデバイス間データ転送」にフォーカスして、そのための研究開発について紹介します。そのための前置きが長くなりますが、ご容赦ください。

鍵は OpenCL BSP の改造

そもそも OpenCL BSP って何？

第2回目の記事で紹介しましたが、FPGAに慣れ親しんでいない人達が参入しやすいように、今や FPGA はハードウェア記述言語だけでなく C や C++ といった一般的なプログラミング言語で利用できるようになりました。そのことを踏まえて、Cygnus には OpenCL が利用できる Intel FPGA ボード Bittware 520N が搭載されています。ユーザーは Intel が提供する OpenCL ベースの FPGA 開発環境である Intel FPGA SDK for OpenCL を使って、FPGA にオフロードするアプリケーションを開発できます。

Intel FPGA SDK for OpenCL でのプログラミングモデルについて復習ですが、ユーザーはホスト PC で動作するホストコードと FPGA で動作するカーネルコードの2種類を記述するだけで、ハードウェア記述言語を知らなくても FPGA を利用できるというものでした。ホストコードは主に OpenCL API (Application Programming Interface) を用いての FPGA のコンフィグレーション、メモリ管理、カーネル実行管理などの FPGA デバイスの制御を担当し、カーネルコードは FPGA にオフロードされる演算を担当します (プログラミングモデルの図は第2回目の記事に掲載しているので、ここでは割愛します)。

重要なのはここからです。下にホストコードとカーネルコードが実際にどこで動いているのかを示しました。gcc のような C コンパイラによって、ホストコードからホストアプリケーションの実行バイナリが生成されます。一方、カーネルコードからは、Intel FPGA SDK for OpenCL に付属している専用コンパイラによって、カーネルコードに記述されている演算をパイプライン処理するハードウェアが生成されます。

Intel FPGA SDK for OpenCL プラットフォームの構成図

赤い破線で囲まれた PCIe コントローラやデバイスドライバ、FPGA ボードに搭載された外部メモリにアクセスするためのコントローラといった (多くの人が見たくない、ぜひ抽象化してほしいと思っているであろう) コンポーネントは、BittWare や Terasic などのFPGA ボードの開発元から提供される OpenCL Board Support Package (BSP) に同梱されています。

当然ですが、FPGA ボードの構成はそれぞれ異なるため、その差異を吸収するために、ボード固有のパラメータや回路は BSP という形で提供されます。そして BSP は、カーネルコードのコンパイル時に自動的に読み込まれ、FPGA の回路情報の生成に利用される仕組みとなっています。

このような仕組みのため、ユーザーは、FPGA を使うにあたって必要な (泥臭い) 部分を意識せずに、ホストコードとカーネルコードの実装に注力すればよく、たとえ異なる FPGA ボードを利用するとしても、その FPGA ボードの BSP が提供されていれば、既存のコードを移植できるという開発モデルになっているわけです。BSP についてもっと知りたい方は、下の Intel の YouTube 動画を参考にしてみて下さい。

BSP Getting Started

ただし、BSP は OpenCL プログラミングを可能にする最低限のインターフェースです。すなわち外部メモリコントローラと PCIe コントローラ、デバイスドライバしか提供していない場合がほとんどです。「FPGA に搭載されているハードウェア (例: Ethernet ポート) を OpenCL から叩きたいな」と思っても、そのようなハードウェアコントローラが BSP に含まれていなければ、その願いを実現させることはできません。

そのような場合には、操作したい外部ペリフェラルを制御するためのハードウェアモジュールをユーザー自身が実装し、それを BSP に適切に組み込む必要があります。つまり、OpenCL カーネルコードから外部ペリフェラルにアクセスできるようにするためには、BSP を改造する必要があるということです。

OpenCL BSP の改造

それでは、どのように BSP を改造するのかについて紹介していきます。BSP を改造して、OpenCL から外部ペリフェラルを叩けるようになるための道のりは、大まかに、次のような 4 つのステップを経ることになります。

  1. 外部ペリフェラルを制御するためのハードウェアモジュールを実装する
  2. BSP に適切に組み込む
  3. OpenCL からハードウェアモジュールにアクセスするための XML を書く
  4. ハードウェアモジュールにアクセスするための OpenCL コードを書く

ここでは、FPGA ボードに搭載されている LED を OpenCL から光らせる事例をベースに説明していきます。その開発フローを図に示すと次のようになります。

FPGA ボードに搭載されている LED を OpenCL から光らせるまでの道のり

この例では、OpenCL カーネルから LED コントローラにデータを送信し、コントローラはそのデータにしたがって LED の発光パターンを制御します。まずは、LED を制御するためのハードウェアモジュールを用意する必要があるので、それを実装します。この例では、Verilog HDL によって実装しました。

次に、実装したコントローラを BSP に組み込みます。この例では、Qsys と呼ばれる Intel FPGA 開発ツール (Xilinx の Vivado IP Integrator に相当するツールです) を利用しました。Qsysのスクリーンショット中の led_st_0 が Verilog HDLで実装された LED コントローラを表します。

BSP にコントローラを組み入れたら、それを OpenCL カーネルからアクセスするための設定が必要です。BSP の hardware ディレクトリの中に、board_spec.xml という FPGA チップおよび外部ペリフェラルのスペックや OpenCL カーネルとの関係性を表す定義ファイルがあるので、図で示すコードをそのファイルに追加します。name、port、type、width、chan_id はそれぞれ、OpenCL カーネルがアクセスするモジュール名、モジュールのポート名、流れるデータの向き、データ幅、OpenCL カーネルがモジュールにアクセスするためのインターフェース名を表します。

そして最後に、LED コントローラにアクセスするための OpenCL カーネルコードを実装します。Intel FPGA SDK for OpenCL では、ユーザーが独自開発したハードウェアモジュールを OpenCL カーネルコードからアクセスするためのベンダー拡張機能である I/O Channel API が提供されています。#pragma ディレクティブによってカーネルコードで I/O Channel API を利用することをコンパイラに明示することによって、実装した LED コントローラにアクセスするための OpenCL カーネルコードを記述できるようになります。

LED コントローラにアクセスするための channel 変数 outLED を宣言し、それを write channel intel() 関数で用いることで、OpenCL カーネルから LED コントローラに対してデータを送信しています。channel 変数を宣言する際 depth と io は、それぞれ OpenCL カーネルと LED コントローラとの間に挿入される FIFO バッファの深さ (エントリ数)、board_spec.xml に追加したコードの chan_id を指定するための attribute です。

このようなステップを経ることによって、OpenCL カーネルから LED を制御できるようになりました。ホストアプリケーションから送信されるデータを OpenCL カーネルが受け取り、そのデータを I/O Channel API によって LED コントローラに送信し、データの値に応じて LED の発光パターンが変化する実験の様子を下に示します。

OpenCL カーネルが LED を制御するのを確認する実験の様子

なお、この実験では、Terasic DE1-SoC 開発キットを使っているので、ホストコードとカーネルコードとの間のインターフェースは PCIe ではないことに注意して下さい。ですが、PCIe 接続のFPGA ボードであってもこれと同じことを実現することができます。

ここでやっと話を戻しますが、我々はこの技術をベースに「計算科学の研究者・開発者がなるべく利用しやすく、かつ高性能なデバイス間データ転送」を実現するための手法を開発しているということです。その具体的な内容をこれから紹介していきます。

実は、「OpenCL から I/O を叩いてみた」という事例は以前からいくつかあるので、新技術というわけではありません。筆者の知る限りでは、一番最初に実現されたのは、下のスライドにあるように2016年ですし、筆者達はここから着想を得ています (みよし様、ありがとうございました)。

また、他のブログや書籍でも紹介されているので、興味のある方はチェックしてみてください。あと、Intelも 独自FPGA ボードのための BSP 作成チュートリアル動画を YouTube にアップロードしています。

Building Custom Platforms for Intel® FPGA SDK for OpenCL™: BSP Basics

OpenCL からキックできる GPU-FPGA 間 DMA 転送

まずは、GPU と FPGA を搭載した計算ノードにおけるデバイス間連携をシームレスに行うための GPU-FPGA 間 DMA データ転送について紹介していきます。

次の図は、我々が開発した GPU-FPGA 間 DMA 転送によるデータ通信と DMA を使わない場合の GPU-FPGA 間データ通信 (従来手法と表記) を表しています。

CPU を経由する GPU-FPGA 間データ通信と DMA を活用したデータ通信

図に示している通り、CPU を経由する間接的な GPU-FPGA 間通信はオーバーヘッドが大きいので、提案手法では CPU を介さないデータ転送 (GPU デバイスメモリへのデータの書き込み・読み出し) を FPGA が自律的に実行します。そして、計算科学の研究者・開発者のようなアプリ屋さんがなるべく利用しやすいように、OpenCL のコードからこの機能を制御できるようにしました。

この提案手法は、BSP に同梱されている PCIe コントローラに GPU デバイスメモリとの DMA の機能を付加し、それを I/O Channel API によって操作することで実現しています。そのため、下のような OpenCL カーネルコードを記述するだけで、GPU-FPGA 間の DMA データ転送を行うことができます。

#pragma OPENCL EXTENSION cl_intel_channels : enable	

typedef struct __attribute__((packed)) cldesc {
  ulong src;
  ulong dst;
  uint  id_and_len;
  uint  unused0;
  uint  unused1;
  uint  unused2;
} cldesc_t;

channel cldesc_t fpga_dma __attribute__((depth(0))) 
     __attribute__((io("chan_fpga_dma")));
channel ulong dma_stat __attribute__((depth(0))) 
     __attribute__((io("chan_dma_stat")));

__kernel void fpga_dma(__global float *restrict fpga_mem,	
                       const ulong gpu_memadr,
                       const uint id_and_len)
{
  cldesc_t desc;
  // DMA transfer GPU -> FPGA	
  desc.src = gpu_memadr;
  desc.dst = (ulong)(&fpga_mem[0]);
  desc.id_and_len = id_and_len;
  write_channel_intel(fpga_dma, desc);
  ulong status = read_channel_intel(dma_stat);
}

このコードは GPU から FPGA への DMA 転送を実行する OpenCL カーネルコードです。1行目の pragma は Intel FPGA SDK for OpenCL の独自拡張である channel の有効化をコンパイラに指示するためのものです。3 ~ 10行目で PCIeコントローラ内に搭載されている DMA コントローラを操作するための構造体を、12、13行目で I/O Channel 変数である fpga_dma と dma_stat を定義しています。生成された構造体は write_channel_intel() 関数によって、PCIe コントローラに送信され、GPUからFPGAへのDMA転送が起動します。そして、DMA の実行が終了したフラグを read_channel_intel() 関数によって読み取っています。

この研究成果は、AsHES 2019 @ リオデジャネイロで発表しました。論文も公開されていますので、詳細を知りたい方は是非ご一読下さい。

OpenCL からキックできる FPGA-FPGA 間通信

次に、複数の FPGA で並列計算を行うための通信技術について紹介していきます。これも BSP の改造を前提とした技術です。具体的には、下の図のように 2D トーラスネットワーク用のハードウェアルータを実装し、BSPに組み入れています (2D トーラスなのは Cygnus のハードウェア構成に依存しています)。この技術によって、ユーザーは OpenCL のレベルで FPGA 間通信かつ通信にタイトに結びついた演算処理を実現するアプリケーションを実装することができます。

ハードウェアルータと OpenCL カーネルの接続関係

下のコードでは、FPGA 1 でデータをセットし、FPGA 2 に送信しています。ユーザーは FPGA 間で通信させたいデータを I/O Channel API によってセットすればよく、I/O Channel から先は、ハードウェアルータが指定した FPGA までデータを転送します。このように、ユーザーは低レイヤーのコンポーネントを意識することなく、FPGA 間通信を利用できるというわけです。

この研究成果は、AsHES 2020 @ Zoom によるオンライン会議で今週の火曜日に発表しました タイトルは Performance Evaluation of Pipelined Communication Combined with Computation in OpenCL Programming on FPGAです。

直近の研究成果なので論文は IEEE Xplore にまだ公開されていないですが、近日中に公開されますので、詳細を知りたい方は是非ご一読下さい。

おまけ: BSP を改造してみたい方へ

本記事では BSP の改造について割と軽く話してしまいましたが、正直言うと中々はまりどころが多くてしんどいです。筆者の同僚が BSP を改造する際に役に立つ Tips をまとめてくれたので、BSP の改造に興味のある方は参考にしてみて下さい。

注意: BSP の改造は完全に自己責任でお願い致します。

まとめ

本記事は、FPGA を導入した国内スパコンである Cygnus をしっかりと使い倒すには様々な課題があることについて触れ、そのうちの 1 つである「計算科学の研究者・開発者がなるべく利用しやすく、かつ高性能なデバイス間データ転送」についての研究事例を紹介しました。

次はいよいよ最終回ですが、そこでは現在取り組んでいる課題や今後の展望について (記事に書ける範囲でですが) 紹介していきたいと思います。どうぞお楽しみに。

筑波大学 計算科学研究センター 小林諒平

タイトルとURLをコピーしました