インテル® HLS コンパイラー用のコードをoneAPI IP Authoring Flow用に書き換える

はじめに

こんにちは。

最近こんな記事を書きました。

この記事を書いた動機の一つに、インテル® HLS コンパイラーのサポートが遠からず終了してしまうというのがあります。そのため、今後新規開発するIPについては是非oneAPI IP Authoring Flowを使って頂きたいわけです。ちなみに、ドキュメントはこちら。

で、新規開発のコードは良いのですが、インテル® HLS コンパイラー対象で書いたコードをoneAPIに持ってくる必要がある場面も出てくるかと思います。この場合、oneAPIなので、IPを開発する場合でもSYCLで記述する必要があり、コードの変更は必須です。

そこで今回は、筆者が個人的に一番応用が効くと思っているHLSコードの移植の仕方の、基本の基本を紹介します。

なお、この記事では、oneAPI Base Toolkitやインテル® HLS コンパイラー、Questa等々のツール類はインストール済みとして話を進めます。まだの方は、上記の記事も参考にしてください（Windows用の記事ですが、Linuxユーザの方にもある程度参考になると思います。というか、Linux版も書いたほうが良いのかな…）。

対象者

この記事は、読者にこれぐらいの知識があることを前提としています。
・インテル® HLS コンパイラーについてはある程度理解し、コードもかける
・C++について、初歩的なこと（クラスがあって、クラスにはメンバがあって、それらの呼び出し方はこうで、等）は理解している
・SYCLについて、大まかな説明は読んだり聞いたりしたことがある（カーネルというのがあって、カーネルとホストは何となくこう書いてつなげる、等）

インテル® HLS コンパイラー用のコードとSYCLコードの違い

ご存じの通り、代表的なインテル® HLS コンパイラー（以下、HLS）用コードは以下のような構成になっています（もちろん他の書き方も出来ますが）。

HLSコード例.cpp

#include "HLS/hls.h"
#include <stdio.h>
#include <stdlib.h>

component void test(...)
{
 // 実際にRTL化されるコード
}

int main()  // テストベンチ記述
{
  // バッファ定義（STLとかmallocとか配列とか）

  // バッファ初期化（テストベクタのセット）

  // component実行
  test(...);

  // 期待値比較

  // 終了処理（バッファ解放等）
}

コメントだけ書いてある箇所は、実際にはコードがあるが省略していると読んでください。

一方、SYCLのコードはこんな感じです。ただし、ここではIP Authoring Flowに向いている
・RTL化する部分（カーネル）はファンクタ（Functor。関数オブジェクトとも）として記述
・カーネルがアクセスするバッファについてはUSMを使う
という記述方法を前提とします。
ファンクタというと難しそうですが、まあクラス定義内にIPの実装を書く、というような理解で大丈夫です。USMとはUnified Shared Memoryの略で、要するに共有メモリですね。IP Authoring Flowにおいては、IP化するモジュールとホスト側との共有メモリ、というようなイメージになります。USM詳細については、こことか、

こことかを参考にしてください。

ともあれ、USMについてはなんとなく理解できていれば大丈夫です。

SYCLコード例.cpp

#include <sycl/ext/intel/fpga_extensions.hpp>
#include <sycl/sycl.hpp>

class Test {
  // 実際にRTL化されるコード
}

int main()  // テストベンチ記述
{
  // selectorの定義
#if FPGA_SIMULATOR
  auto selector = sycl::ext::intel::fpga_simulator_selector_v;
#elif FPGA_HARDWARE
  auto selector = sycl::ext::intel::fpga_selector_v;
#else  // #if FPGA_EMULATOR
  auto selector = sycl::ext::intel::fpga_emulator_selector_v;
#endif

  // queueの作成
  sycl::queue q(selector);

  // バッファ定義（USM使用）
  int *a = sycl::malloc_shared<int>(1000, q);

  // バッファ初期化（テストベクタのセット）

  // カーネル（クラス）実行
  q.single_task(Test{a, 1000}).wait();

  // 期待値比較

  // 終了処理（バッファ解放等）
  sycl::free(a, q);

}

こちらも、コメントだけ書いてある箇所は、実際にはコードがあるが省略していると読んでください。ちょっとだけ実装が書いてある部分もありますが、これについても一部の実装は省略されています。

HLSコードとの大きな違いは、実際にRTL化するモジュールをclassの中に書くということと、SYCLのお約束の記述をmain()内にしないといけないということです。具体的には、

Selectorを定義する

この部分です。

SYCLコード例.cpp

#if FPGA_SIMULATOR
  auto selector = sycl::ext::intel::fpga_simulator_selector_v;
#elif FPGA_HARDWARE
  auto selector = sycl::ext::intel::fpga_selector_v;
#else  // #if FPGA_EMULATOR
  auto selector = sycl::ext::intel::fpga_emulator_selector_v;
#endif

Selectorとは、Deviceとも呼ばれるアクセラレータ（以下、SYCL用語に合わせてDeviceと書きます）を選択するためのオブジェクトで、ここではDeviceとして明示的にFPGAを選択するので、FPGA selectorというのを使います。
ただ、FPGAの場合、Selectorが３種類あるんですよね。具体的にはEmulation用のSelector、Simulation用のSelector、実機動作用のSelectorです。ですので、通常は上記のようにifdefを使って記述し、コンパイル時にどのSelectorを使うか選べるようにするのが一般的です。
ちなみに、IP Authoring Flowでは実機動作を行うことはあまりないので、EmulationおよびSimulationのSelectorだけでもOKです。
あと、EmulationとSimulationの違いですが、EmulationはあくまでCPUを使ってエミュレーションするもの、Simulationは実際にRTLを合成してQuesta等のHW Simulatorでシミュレーションするものです。

ということで、IP Authoring Flow時は、おまじない的に上記をそのままコピペで記述してしまうのでOKです。

Queueを定義する

この部分です。

SYCLコード例.cpp

  sycl::queue q(selector);

SYCLにおけるQueueとは、使用するDeviceに対する仕事を送り込むためのキューというような意味合いで、Device毎に独立したQueueが必要なため、インスタンス時にはSelectorを指定します。
これも、IP Authoring Flow時は、おまじない的に上記をそのままコピペで記述しまえばほぼOKです。

USMのバッファを定義する

この部分です。

SYCLコード例.cpp

  int *a = sycl::malloc_shared<int>(1000, q);

通常は、このsycl::malloc_sharedというAPIを使います。上記で定義したキューが引数に指定されていることに注意してください。例では共有メモリが１つだけ定義されていますが、実際は複数定義されることが多いです。
USMは前述の通り共有メモリなので、Hostから初期値（テストベクタ）を書き込み、Deviceでそれを読み込んで動作させ結果をまた共有メモリに書き込み、その後Hostがそれを読み込んで期待値比較する、というようなことができます。

この記述だけはコピペでは無理なので、必要なバッファ数・サイズを考えて記述する必要があります。

カーネルを実行する

この部分です。

SYCLコード例.cpp

  q.single_task(Test{a, 1000}).wait();

またキューが出てきましたね。先ほど共有メモリの確保が既にキューに積まれました。ここでさらに、カーネルをキューに積み、これによりカーネルの実行が開始されます。.wait()という記述がついているので、カーネル実行終了までこの行はブロックされます。
single_taskという関数は文字通りキューにシングルの（並列ではない）タスクを積むということで、IP Authoring Flowでは常にこのsingle_taskを使います。

いかがでしょうか。IP Authoring Flowにおいては、SYCLとはいえ記述することは大体いつも同じなので、HLSコードの内容をベースとなるSYCLコードにパカパカ当てはめていけば比較的簡単にSYCL化出来そうな気がしませんか？
では、次から実際に移植作業を見ていきましょう。

移植元のHLSコード

今回移植対象とするのは、インテル® HLS コンパイラーをインストールするとexampleディレクトリに展開される以下のコードとします。なぜこれを選んだかというと、メモリI/Fを詳細定義（バス幅やバックプレッシャー有無等）しているという点で実戦向きと考えたためです。
以下にソースコードを掲載します。これは、インテル® HLS コンパイラーをインストールするとhls/examples/tutorials/interfaces/pointer_mm_hostに展開されているはずです。内容としては単なるVector Additionですが、mm_hostを使って３系統のメモリバスを定義していることに注意してください。

part_3_ddr_hosts.cpp

//  Copyright (c) 2023 Intel Corporation                                  
//  SPDX-License-Identifier: MIT                                          


#include "HLS/hls.h"

#include <stdio.h>
#include <stdlib.h>

component void vector_add(ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& a, // bank 1
                          ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& b, // bank 1
                          ihc::mm_host<int, ihc::aspace<2>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& c, // bank 2
                          int N) {
  for (int i = 0; i < N; ++i) {
    c[i] = a[i] + b[i];  
  }
}

#define TEST_SIZE 1000
#define SEED 4

int main(void) {

  int A[TEST_SIZE];
  int B[TEST_SIZE];
  int C[TEST_SIZE];

  // mm_host interface class instances
  ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> > mm_A(A, sizeof(int)*TEST_SIZE);
  ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> > mm_B(B, sizeof(int)*TEST_SIZE);
  ihc::mm_host<int, ihc::aspace<2>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> > mm_C(C, sizeof(int)*TEST_SIZE);

  // prepare the input data
  srand(SEED);
  for (int i = 0; i < TEST_SIZE; ++i) {
    A[i] = rand();
    B[i] = rand();
  }

  // Run the component
  vector_add(mm_A, mm_B, mm_C, TEST_SIZE);
  
  // Check the output
  bool passed = true;
  for (int i = 0; i < TEST_SIZE; ++i) {
    bool data_okay = (C[i] == (A[i] + B[i]));
  
    passed &= data_okay;
    if (!data_okay) {
      printf("ERROR: C[%d] = %d != %d\n", i, C[i], (A[i] + B[i]));
    }
  }

  if (passed) {
    printf("PASSED\n");
  } else {
    printf("FAILED\n");
  }

  return 0;
}

ベースとするSYCLコード

上記のHLSコードを、以下のoneAPIのサンプルにはめ込んでいってSYCL化します。このサンプルは、IP Authoring Flowを使う上で最も基本となる（と筆者が考えている）もので、前述のファンクタとUSMが使われています。

ちなみに、上記のURLはmasterブランチですが、適宜読み替えてください。例えば、現時点でのoneAPI Base Toolkitの最新版は2023.2ですが、その場合は以下を見るのがおすすめです。これ以降は、このtags/2023.2.0を見るようにします。

ソースコードについては、以下のvector_add.cppをリネームして使っていきましょう。

移植開始！

ではまず、このSYCLのサンプルを適当な場所にコピーします。筆者は、ルートにworkというディレクトリを掘って、そこにコピーしました。

$ git clone https://github.com/oneapi-src/oneAPI-samples
$ cd oneAPI-samples
$ git checkout tags/2023.2.0
$ cp -r DirectProgramming/C++SYCL_FPGA/Tutorials/GettingStarted/fpga_compile/part2_dpcpp_functor_usm ~/work

vector_addという名前のまま作業しても良いんですが、せっかくですので名前を変えてみます（移植元のHLSコード及び移植先のoneAPIコードどちらもvector addなので、変える必要ないといえばないんですが…）。ここではtestという何のひねりもない名前にしてみます。

$ cd work/part2_dpcpp_functor_usm/src
$ mv vector_add.cpp test.cpp

srcディレクトリにCMakeLists.txtがあります。vector_addという名前を変更したので、これをちょっと修正する必要があります。

CMakeLists.txt（修正前）

set(SOURCE_FILE vector_add.cpp)
set(TARGET_NAME vector_add)
set(EMULATOR_TARGET ${TARGET_NAME}.fpga_emu)
set(SIMULATOR_TARGET ${TARGET_NAME}.fpga_sim)

上の２行のvector_addをtestに変更しておきます。

CMakeLists.txt（修正後）

set(SOURCE_FILE test.cpp)
set(TARGET_NAME test)
set(EMULATOR_TARGET ${TARGET_NAME}.fpga_emu)
set(SIMULATOR_TARGET ${TARGET_NAME}.fpga_sim)

クラス（ファンクタ）記述の修正

さて、まずクラスの定義部分を修正していきましょう。元々はこんな記述です。IPの入出力がメンバ変数として定義され、実際のIPの動作部分はoperatorとして記述されています。なぜclassではなくstructかというと、多分そのほうが楽なためです（笑）。実際、どちらを使ってもOKですが、個人的にはIP Authoring Flowではstructで十分かなと思います。

test.cpp

struct VectorAdd {
  int *const vec_a_in;
  int *const vec_b_in;
  int *const vec_c_out;
  int len;

  void operator()() const {
    for (int idx = 0; idx < len; idx++) {
      int a_val = vec_a_in[idx];
      int b_val = vec_b_in[idx];
      int sum = a_val + b_val;
      vec_c_out[idx] = sum;
    }
  }
};

また、移植したいHLSのcomponent記述（RTL化する記述・関数）は以下です。

part_3_ddr_hosts.cpp

component void vector_add(ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& a, // bank 1
                          ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& b, // bank 1
                          ihc::mm_host<int, ihc::aspace<2>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& c, // bank 2
                          int N) {
  for (int i = 0; i < N; ++i) {
    c[i] = a[i] + b[i];  
  }
}

すみません、どちらもvector additionなのでちょっと分かりづらいですが…
ここでは、オリジナルのHLSコードからcomponent記述をそのままコピペでSYCL側に持ってくることを想定します。そうすると、イメージ的にはクラスの記述はこんな感じになりますかね。

test.cpp

// 実際はSYCLでihc::mm_hostは使えないので、このソースコードは無効！
void vector_add(ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& a, // bank 1
                ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& b, // bank 1
                ihc::mm_host<int, ihc::aspace<2>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& c, // bank 2
                int N) {
  for (int i = 0; i < N; ++i) {
    c[i] = a[i] + b[i];  
  }
}

struct Test {
  ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> > a;
  ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> > b;
  ihc::mm_host<int, ihc::aspace<2>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> > c;
  int N;

  void operator()() const {
    vector_add(a, b, c, N);
  }
};

ですが、SYCLではihc::mm_hostなんていう定義はないので、ここを書き換える必要があります。具体的には、以下のドキュメントにあるような記述にします。

上記の記述にある通り、まず先頭のinclude周りをこんな感じに変更します。３行ほど追加する感じですね。prototypeとかexperimentalとかいう単語が気になりますが、気にせず行きましょう！

test.cpp

#include <sycl/ext/intel/fpga_extensions.hpp>
#include <sycl/ext/intel/prototype/interfaces.hpp> // 追加
#include <sycl/sycl.hpp>

using namespace sycl; // 追加
using ext::intel::experimental::property::usm::buffer_location; // 追加

さて、ここではihc::mm_hostに変えてregister_map_mmhostを使います。
これが↓、

part_3_ddr_hosts.cpp

ihc::mm_host<int, ihc::aspace<1>, ihc::awidth<32>, ihc::dwidth<256>, ihc::latency<0>, ihc::maxburst<8>, ihc::waitrequest<true> >& a, // bank 1

こう修正されます↓。

test.cpp

  register_map_mmhost(
		      BL0,     // buffer_location or aspace
		      32,      // address width
		      256,     // data width
		      0,       // ! latency, must be atleast 16
		      1,       // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
		      8,       // maxburst
		      0,       // align, 0 defaults to alignment of the type
		      1        // waitrequest, 0: false, 1: true
		      ) int * const a;

上記で、BL0というのは、I/FごとにつけるIDのようなものです。それ以外のバス幅やアドレス幅等、オリジナルのHLSの定義と同じになっていることが分かると思います。

変更後の上記のコード（クラス定義）はこんな感じ。

test.cpp

void vector_add(int* a, int* b, int* c, int N) {
  for (int i = 0; i < N; ++i) {
    c[i] = a[i] + b[i];  
  }
}

static constexpr int BL0 = 0;
static constexpr int BL1 = 1;
static constexpr int BL2 = 2;

struct Test {
  register_map_mmhost(
		      BL0,     // buffer_location or aspace
		      32,      // address width
		      256,     // data width
		      0,       // ! latency, must be atleast 16
		      1,       // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
		      8,       // maxburst
		      0,       // align, 0 defaults to alignment of the type
		      1        // waitrequest, 0: false, 1: true
		      ) int * const a;
  register_map_mmhost(
		      BL1,     // buffer_location or aspace
		      32,      // address width
		      256,     // data width
		      0,       // ! latency, must be atleast 16
		      1,       // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
		      8,       // maxburst
		      0,       // align, 0 defaults to alignment of the type
		      1        // waitrequest, 0: false, 1: true
		      ) int * const b;
  register_map_mmhost(
		      BL2,     // buffer_location or aspace
		      32,      // address width
		      256,     // data width
		      0,       // ! latency, must be atleast 16
		      2,       // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
		      8,       // maxburst
		      0,       // align, 0 defaults to alignment of the type
		      1        // waitrequest, 0: false, 1: true
		      ) int * const c;
  int N;

  void operator()() const {
    vector_add(a, b, c, N);
  }
};

注目してほしいのは、少なくともここでは、HLSコードから持ってきたRTL化したい関数の内容はI/F以外修正の必要が無かったです。つまり、SYCLをIP Authoring Flow用に使うと割り切れば、C++の知識もそれほど必要ではなく、Cの範囲内で戦えます（異論はあるかと思いますが）。

テストベンチ（main関数）の修正

次に、カーネル以外の部分を修正していきます。IP Authoring Flowなので、テストベンチ・テストベクタに当たる部分の修正ですね。
まずバッファ周りを修正します。前述の通り、バッファはUSMを使います。
実際のコードでは、具体的には以下の部分

test.cpp

    // declare arrays and fill them
    // allocate in shared memory so the kernel can see them
    int *vec_a = sycl::malloc_shared<int>(kVectSize, q);
    int *vec_b = sycl::malloc_shared<int>(kVectSize, q);
    int *vec_c = sycl::malloc_shared<int>(kVectSize, q);

ここがUSMの記述ですね。これを、元々のHLSコードに合わせてこんな感じに修正してみます。

test.cpp

    int *A = sycl::malloc_shared<int>(sizeof(int)*TEST_SIZE, q, property_list{buffer_location(BL0)});
    int *B = sycl::malloc_shared<int>(sizeof(int)*TEST_SIZE, q, property_list{buffer_location(BL1)});
    int *C = sycl::malloc_shared<int>(sizeof(int)*TEST_SIZE, q, property_list{buffer_location(BL2)});

property_list{buffer_location(BL0)}というような引数が最後のところに追加されていることに注意してください。これは、ここで定義したバッファが上記でregister_map_mmhostで定義した３つのメモリI/Fのうちのどれに接続されるかを示すものと考えてください。
USMと言っても普通の配列のようにアクセスできますので、データの初期化はオリジナルの記述そのままで良さそうです。

カーネルのサブミットのところはこんな感じに書き換えます。

test.cpp

    q.single_task<VectorAddID>(Test{A, B, C, TEST_SIZE})
        .wait();

<VectorAddID>のところはSYCL上はあまり意味がないんですが、後でレポートを見る際に見やすくなる効果があります。ここは後ほどTestIDみたいな名前に修正します。
期待値比較のところは、オリジナルのHLSのコードそのままで良さそうです。
最終的に、コードはこんな感じになりました。

test.cpp

#include <iostream>

// oneAPI headers
#include <sycl/ext/intel/fpga_extensions.hpp>
#include <sycl/ext/intel/prototype/interfaces.hpp>
#include <sycl/sycl.hpp>

using namespace sycl;
using ext::intel::experimental::property::usm::buffer_location;

// Forward declare the kernel name in the global scope. This is an FPGA best
// practice that reduces name mangling in the optimization reports.
class TestID;

void vector_add(int* a, int* b, int* c, int N) {
  for (int i = 0; i < N; ++i) {
    c[i] = a[i] + b[i];  
  }
}

static constexpr int BL0 = 0;
static constexpr int BL1 = 1;
static constexpr int BL2 = 2;

struct Test {
  register_map_mmhost(
		      BL0,     // buffer_location or aspace
		      32,      // address width
		      256,     // data width
		      0,       // ! latency, must be atleast 16
		      1,       // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
		      8,       // maxburst
		      0,       // align, 0 defaults to alignment of the type
		      1        // waitrequest, 0: false, 1: true
		      ) int * const a;
  register_map_mmhost(
		      BL1,     // buffer_location or aspace
		      32,      // address width
		      256,     // data width
		      0,       // ! latency, must be atleast 16
		      1,       // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
		      8,       // maxburst
		      0,       // align, 0 defaults to alignment of the type
		      1        // waitrequest, 0: false, 1: true
		      ) int * const b;
  register_map_mmhost(
		      BL2,     // buffer_location or aspace
		      32,      // address width
		      256,     // data width
		      0,       // ! latency, must be atleast 16
		      2,       // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
		      8,       // maxburst
		      0,       // align, 0 defaults to alignment of the type
		      1        // waitrequest, 0: false, 1: true
		      ) int * const c;
  int N;

  void operator()() const {
    vector_add(a, b, c, N);
  }
};

#define TEST_SIZE 1000
#define SEED 4

int main() {
  bool passed = true;
  try {
    // Use compile-time macros to select either:
    //  - the FPGA emulator device (CPU emulation of the FPGA)
    //  - the FPGA device (a real FPGA)
    //  - the simulator device
#if FPGA_SIMULATOR
    auto selector = sycl::ext::intel::fpga_simulator_selector_v;
#elif FPGA_HARDWARE
    auto selector = sycl::ext::intel::fpga_selector_v;
#else  // #if FPGA_EMULATOR
    auto selector = sycl::ext::intel::fpga_emulator_selector_v;
#endif

    // create the device queue
    sycl::queue q(selector);

    auto device = q.get_device();

    std::cout << "Running on device: "
              << device.get_info<sycl::info::device::name>().c_str()
              << std::endl;

    if (!device.has(sycl::aspect::usm_host_allocations)) {
      std::terminate();
    }

    // declare arrays and fill them
    // allocate in shared memory so the kernel can see them
    int *A = sycl::malloc_shared<int>(sizeof(int)*TEST_SIZE, q, property_list{buffer_location(BL0)});
    int *B = sycl::malloc_shared<int>(sizeof(int)*TEST_SIZE, q, property_list{buffer_location(BL1)});
    int *C = sycl::malloc_shared<int>(sizeof(int)*TEST_SIZE, q, property_list{buffer_location(BL2)});

    // prepare the input data
    srand(SEED);
    for (int i = 0; i < TEST_SIZE; ++i) {
      A[i] = rand();
      B[i] = rand();
    }

    q.single_task<TestID>(Test{A, B, C, TEST_SIZE})
        .wait();

    // Check the output
    bool passed = true;
    for (int i = 0; i < TEST_SIZE; ++i) {
      bool data_okay = (C[i] == (A[i] + B[i]));
      
      passed &= data_okay;
      if (!data_okay) {
	printf("ERROR: C[%d] = %d != %d\n", i, C[i], (A[i] + B[i]));
      }
    }
    
    if (passed) {
      printf("PASSED\n");
    } else {
      printf("FAILED\n");
    }
    
    sycl::free(A, q);
    sycl::free(B, q);
    sycl::free(C, q);
  } catch (sycl::exception const &e) {
    // Catches exceptions in the host code.
    std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n";

    // Most likely the runtime couldn't find FPGA hardware!
    if (e.code().value() == CL_DEVICE_NOT_FOUND) {
      std::cerr << "If you are targeting an FPGA, please ensure that your "
                   "system has a correctly configured FPGA board.\n";
      std::cerr << "Run sys_check in the oneAPI root directory to verify.\n";
      std::cerr << "If you are targeting the FPGA emulator, compile with "
                   "-DFPGA_EMULATOR.\n";
    }
    std::terminate();
  }
  return passed ? EXIT_SUCCESS : EXIT_FAILURE;
}

いかがでしょう。修正箇所はそれほど多くなかったと思いません？

実行してみる

では、いつもの通り実行してみます。ここでは、対象デバイスをAgilex7としました。まずはEmulationから。

$ pwd
/home/demo/work/part2_dpcpp_functor_usm
$ mkdir build
$ cd build
$ cmake -DFPGA_DEVICE=Agilex7 ..
(snip)
$ make fpga_emu
(snip)
$ ./test.fpga_emu 
Running on device: Intel(R) FPGA Emulation Device
PASSED

make fpga_emuは単にmakeだけでも大丈夫のはずです。上記のように動きましたか？

では、本命のSimulationです。以下のように実行します。

$ make fpga_sim
(snip)
$ export CL_CONTEXT_MPSIM_DEVICE_INTELFPGA=1
$ ./test.fpga_sim
Running on device: SimulatorDevice : Multi-process Simulator (aclmsim0)
PASSED
$ unset CL_CONTEXT_MPSIM_DEVICE_INTELFPGA

うまく動作しましたか？
ベースにoneAPIの公式サンプルを使ったため、makeもそのまま使えて便利ですね！

最後に

たまたま複数方面から、HLSコードをoneAPIに持っていくやり方が分からんという相談を受けたので、慌ててこんな記事を書いてみました。
次回は、実際に出来上がったIPのRTL上のI/Fを見ていくような記事を書きたいと思っています。まだoneAPIのIP Authoring Flowは出来たばかりで、実際に生成されるIPのバス名称が分かりにくかったり、必要のないバスが生成されたりすることがあるので。

ではまた。

Notices & Disclaimers

Intel technologies may require enabled hardware, software or service activation.

No product or component can be absolutely secure.

Your costs and results may vary.

© Intel Corporation. Intel, the Intel logo, and other Intel marks are trademarks of Intel Corporation or its subsidiaries. Other names and brands may be claimed as the property of others.