ホスト アプリケーション

Vitis コア開発キットでは、ホスト コードは業界標準の OpenCL API を使用した C または C++ 言語で記述します。Vitis コア開発キットには、OpenCL 1.2 エンベデッド プロファイル準拠ランタイム API が含まれます。

通常、ホスト コードは次の 3 つのセクションに分けることができます。

  1. 環境の設定。
  2. 1 つまたは複数のカーネルの実行を含むコア コマンドの実行。
  3. リソースのポスト プロセスとリリース。
注記: Vitis コア開発キットでは、OpenCL インストーラブル クライアント ドライバー (ICD) 拡張 (cl_khr_icd) がサポートされます。これにより、OpenCL を複数インプリメンテーションして、同じシステム内に共存させることができます。インストール方法の詳細は、OpenCL インストーラブル クライアント ドライバー ローダー を参照してください。
注記: ホスト プログラムをマルチスレッドで実行する場合は、Vitis コア開発キットアプリケーションからの fork() システム呼び出しを呼び出す際に注意が必要です。fork() は、すべてのランタイム スレッドを複製するわけではありません。このため、子プロセスを完全な Vitis コア開発キットのアプリケーションとして実行することはできません。Vitis ソフトウェア プラットフォーム アプリケーションから別のプロセスを起動する場合は、posix_spawn() を使用することをお勧めします。

OpenCL 環境の設定

Vitis コア開発キットのホスト コードは、OpenCL プログラミングの基本枠組みに従います。この環境を正しく設定するには、ホスト アプリケーションがターゲット プラットフォーム、デバイス、コンテキスト、コマンド キュー、プログラムなどの標準 OpenCL 構造を初期化する必要があります。

ヒント: この資料で使用されるホスト コード例および API コマンドは、OpenCL C API に従っています。ただし、XRT では OpenCL C++ ラッパー API もサポートされており、Vitis サンプルの多くが C++ API を使用して記述されています。この C++ ラッパー API の詳細は、https://www.khronos.org/registry/OpenCL/specs/opencl-cplusplus-1.2.pdf を参照してください。

プラットフォーム

初期化の際は、ホスト アプリケーションは 1 つまたは複数のザイリンクス デバイスを含むプラットフォームを特定する必要があります。次のコード部分は、ザイリンクス プラットフォームを特定する共通の方法を示しています。
cl_platform_id platform_id;         // platform id

err = clGetPlatformIDs(16, platforms, &platform_count);
    
// Find Xilinx Platform
for (unsigned int iplat=0; iplat<platform_count; iplat++) {
  err = clGetPlatformInfo(platforms[iplat], 
    CL_PLATFORM_VENDOR, 
    1000, 
    (void *)cl_platform_vendor,
    NULL);

  if (strcmp(cl_platform_vendor, "Xilinx") == 0) { 
  // Xilinx Platform found
  platform_id = platforms[iplat];
  }
}

OpenCL API 呼び出し clGetPlatformIDs を使用して、システムで使用可能な OpenCL プラットフォームを検索します。その後、clGetPlatformInfo を使用して、cl_platform_vendor が文字列 "Xilinx"ザイリンクス デバイス ベースのプラットフォームを取得します。

注記: 上記のコードまたはこの章で使用されているほかのホスト コード例でははっきりと示されていませんが、各 OpenCL API 呼び出しの後にエラー チェックを使用することをお勧めします。そのようにするとデバッグしやすくなり、エミュレーション フローまたはハードウェア実行中にホスト コードおよびカーネル コードをデバッグする場合に生産性が向上します。次に、clGetPlatformIDs コマンドのエラー チェック のコード例を示します。
err = clGetPlatformIDs(16, platforms, &platform_count);
if (err != CL_SUCCESS) {
  printf("Error: Failed to find an OpenCL platform!\n");
  printf("Test failed\n");
  exit(1);
}

デバイス

ザイリンクス プラットフォームが見つかったら、アプリケーションが該当するザイリンクス デバイスを識別する必要があります。

次のコード例では、上限 16 の clGetDeviceIDs API を使用するザイリンクス デバイスすべてを検出しています。
cl_device_id devices[16];  // compute device id
char cl_device_name[1001];
    
err = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ACCELERATOR, 
  16, devices, &num_devices);
    
printf("INFO: Found %d devices\n", num_devices);
    
//iterate all devices to select the target device.
for (uint i=0; i<num_devices; i++) {
  err = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, 1024, cl_device_name, 0);
  printf("CL_DEVICE_NAME %s\n", cl_device_name);
}
重要: clGetDeviceIDs API は platform_id および CL_DEVICE_TYPE_ACCELERATOR を指定して呼び出されており、使用可能なすべてのザイリンクス デバイスが返されます。

サブデバイス

Vitis コア開発キットでは、デバイスに複数の同じカーネルまたは異なるカーネルのインスタンスが含まれることがあります。clCreateSubDevices という OpenCL API を使用してもホスト コードでデバイスを複数のサブデバイスに分割できますが、Vitis コア開発キットでは CL_DEVICE_PARTITION_EQUALLY による均等なサブデバイス分割 (サブデバイスごとにカーネル インスタンスを 1 つずつ含む) がサポートされています。

この後の例では、次が実行されます。

  1. サブデバイスごとに 1 つのカーネル インスタンスを実行する同等分割で作成されます。
  2. 別のコンテキストおよびコマンド キューを使用してサブデバイス リストを反復実行し、それぞれでカーネルを実行します。
  3. 単純にするため、カーネル実行に関する API (および対応するバッファーに関する) コードは示していませんが、関数 run_cu 内に記述してください。
cl_uint num_devices = 0;
  cl_device_partition_property props[3] = {CL_DEVICE_PARTITION_EQUALLY,1,0};
  
  // Get the number of sub-devices
  clCreateSubDevices(device,props,0,nullptr,&num_devices);  
  
  // Container to hold the sub-devices
  std::vector<cl_device_id> devices(num_devices);  

  // Second call of clCreateSubDevices    
  // We get sub-device handles in devices.data()
  clCreateSubDevices(device,props,num_devices,devices.data(),nullptr); 

  // Iterating over sub-devices
  std::for_each(devices.begin(),devices.end(),[kernel](cl_device_id sdev) {
      
	  // Context for sub-device
      auto context = clCreateContext(0,1,&sdev,nullptr,nullptr,&err);  
      
	  // Command-queue for sub-device
      auto queue = clCreateCommandQueue(context,sdev,
      CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,&err); 
      
      // Execute the kernel on the sub-device using local context and 
	queue run_cu(context,queue,kernel); // Function not shown 
  });
重要: 上記の例に示すように、サブデバイスごとに別のコンテキストを作成する必要があります。OpenCL では、複数のデバイスおよびサブデバイスを保持するコンテキストが作成できますが、XRT ではデバイスおよびサブデバイスごとに別のコンテキストを作成する必要があります。

コンテキスト

clCreateContext API を使用して、ホスト マシンと通信する 1 つのザイリンクス デバイスを含むコンテキストを作成します。

context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);

このコード例では、clCreateContext API を使用して、1 つのザイリンクス デバイスを含むコンテキストを作成しています。ザイリンクスでは、デバイスまたはサブデバイスごとに 1 つのコンテキストのみを作成することをお勧めしています。ただし、サブデバイスごとに 1 つのコンテキストが使用されている場合は、ホスト プログラムで複数のコンテキストを使用する必要があります。

コマンド キュー

clCreateCommandQueue API は、デバイスごとに 1 つまたは複数のコマンド キューを作成します。FPGA には、同じカーネルまたは異なるカーネルのいずれかの複数のカーネルを含めることができます。ホスト アプリケーションを開発する場合、デバイスでカーネルを実行するためのプログラミング方法が主に 2 つあります。

  1. 1 つの順不同コマンド キュー: 複数のカーネル実行を同じコマンド キューで要求できます。XRT によりこれらのカーネルができるだけ早く任意の順序で実行され、FPGA 上でのカーネルの同時実行を可能にします。
  2. 複数の順序どおりのコマンド キュー: 各カーネル実行は異なる順序どおりのコマンド キューから要求されます。このような場合、XRT は異なるコマンド キューからカーネルを実行し、それらがデバイス上で同時に実行されるようになるので、パフォーマンスが改善されます。

次に、順不同および順序どおりのコマンド キューを作成する標準 API 呼び出しの例を示します。

// Out-of-order Command queue
commands = clCreateCommandQueue(context, device_id, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);

// In-order Command Queue
commands = clCreateCommandQueue(context, device_id, 0, &err);

プログラム

ビルド プロセス で説明するように、ホストおよびとカーネル コードは別々にコンパイルされ、ホスト プログラム実行ファイルおよび FPGA バイナリ (.xclbin) などの実行ファイルが別々に作成されます。ホスト アプリケーションが実行されると、clCreateProgramWithBinary API を使用して .xclbin を読み込む必要があります。

次のコード例は、標準の OpenCL API を使用して .xclbin ファイルからプログラムをビルドする方法を示しています。

unsigned char *kernelbinary;
char *xclbin = argv[1];

printf("INFO: loading xclbin %s\n", xclbin);
 
int size=load_file_to_memory(xclbin, (char **) &kernelbinary);
size_t size_var = size; 

cl_program program = clCreateProgramWithBinary(context, 1, &device_id, 
					 &size_var,(const unsigned char **) &kernelbinary, 
					 &status, &err);

// Function 
int load_file_to_memory(const char *filename, char **result)
{
  uint size = 0;
  FILE *f = fopen(filename, "rb");
  if (f == NULL) {
    *result = NULL;
    return -1; // -1 means file opening fail
  }
  fseek(f, 0, SEEK_END);
  size = ftell(f);
  fseek(f, 0, SEEK_SET);
  *result = (char *)malloc(size+1);
  if (size != fread(*result, sizeof(char), size, f)) {
    free(*result);
    return -2; // -2 means file reading fail
  }
  fclose(f);
  (*result)[size] = 0;
  return size;
}

このコード例では、次が実行されます。

  1. カーネルのバイナリ ファイル (.xclbin) をコマンド ライン引数 argv[1] から渡します。
    ヒント: コマンド ライン引数から .xclbin を渡すのも 1 つの方法ですが、ホスト プログラムでカーネル バイナリ ファイルをハードコード化したり、環境変数を使用して定義したり、カスタム初期化ファイルから読み込んだり、その他の適切な方法を使用したりもできます。
  2. load_file_to_memory 関数を使用して、ファイルの内容をホスト マシンのメモリ空間に読み込みます。
  3. clCreateProgramWithBinary API を使用して、指定したコンテキストおよびデバイスでプログラム作成プロセスを完了します。

FPGA でのコマンドの実行

OpenCL 環境が初期化されたら、ホスト アプリケーションがデバイスに対してコマンドを発行し、カーネルと対話できるようになります。これらのコマンドには、次が含まれます。

  1. カーネルの設定。
  2. FPGA デバイスのバッファー転送。
  3. FPGA 上でのカーネルの実行。
  4. イベントの同期化。

カーネルの設定

デバイスの識別、コンテキスト、コマンド キュー、およびプログラムの作成などの OpenCL 環境を設定したら、ホスト アプリケーションがデバイスで実行し、カーネルを設定するカーネルを見つける必要があります。

.xclbin ファイル内に含まれるカーネル (program) にアクセスするためには、OpenCL API の clCreateKernel を使用する必要があります。cl_kernel オブジェクトは、FPGA に読み込まれるホスト アプリケーションで実行可能なプログラムのカーネルを識別します。次のコード例では、読み込んだプログラムで定義された 2 つのカーネルを見つけています。

kernel1 = clCreateKernel(program, "<kernel_name_1>", &err);            
kernel2 = clCreateKernel(program, "<kernel_name_2>", &err);  // etc

カーネル引数の設定

Vitis ソフトウェア プラットフォームワークでは、cl_kernel オブジェクトに次の 2 つのタイプの引数を設定できます。

  1. スカラー引数は、定数またはコンフィギュレーション タイプのデータなどの小型のデータに使用します。これらは、ホスト アプリケーションの観点からは書き込み専用、つまりカーネルへの入力です。
  2. バッファー引数は、大型のデータ転送に使用します。値は、プログラムおよびカーネル オブジェクトに関連付けられたコンテキストを使用して作成されたメモリ オブジェクトへのポインターで、カーネルへの入力またはカーネルからの出力になります。
カーネル引数は、次に示すように clSetKernelArg コマンドを使用して設定します。次の例では、2 つのスカラーと 2 つのバッファー引数のカーネル引数を設定しています。
// Create memory buffers
cl_mem dev_buf1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, size, &host_mem_ptr1, NULL);
cl_mem dev_buf2 = clCreateBuffer(context, CL_MEM_READ_ONLY, size, &host_mem_ptr2, NULL);

int err = 0;
// Setup scalar arguments
cl_uint scalar_arg_image_width = 3840;
err |= clSetKernelArg(kernel, 0, sizeof(cl_uint), &scalar_arg_image_width); 
cl_uint scalar_arg_image_height = 2160; 
err |= clSetKernelArg(kernel, 1, sizeof(cl_uint), &scalar_arg_image_height); 
    
// Setup buffer arguments
err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &dev_buf1);
err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &dev_buf2);
重要: OpenCL では、カーネルのエンキュー前であればいつでもカーネル引数を設定できますが、できるだけ早期に設定することをお勧めします。XRT では、デバイスのどこにバッファーを置くかわかる前にバッファーを移行しようとすると、エラー メッセージが表示されます。このため、カーネル引数はエンキュー (たとえば clEnqueueMigrateMemObjects) を実行する前にバッファーに設定するようにしてください。

FPGA デバイスのバッファー転送

ホスト プログラムとハードウェア カーネル間の送受信方法は、デバイス内のグローバル メモリを転送されるデータによって異なります。FPGA とデータの送受信を実行するには、clCreateBufferclEnqueueWriteBuffer および clEnqueueReadBuffer コマンドを使用します。

注記: ザイリンクスでは、clEnqueueReadBuffer および clEnqueueWriteBuffer ではなく、clEnqueueMigrateMemObjects を使用することをお勧めしています。

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

int host_mem_ptr[MAX_LENGTH]; // host memory for input vector
// Fill the memory input
for(int i=0; i<MAX_LENGTH; i++) {
  host_mem_ptr[i] = <... >   
}
cl_mem dev_mem_ptr = clCreateBuffer(context,  CL_MEM_READ_WRITE, 
                     sizeof(int) * number_of_words, NULL, NULL);

err = clEnqueueWriteBuffer(commands, dev_mem_ptr, CL_TRUE, 0,
      sizeof(int) * number_of_words, host_mem_ptr, 0, NULL, NULL);
重要: 現時点では、ホストからグローバル メモリまでのスループットを最大化するために 4 GB を上回るバッファーは使用できませんが、ザイリンクス ではバッファー サイズをできれば少なくとも 2 MB に抑えることをお勧めしています。

単純なアプリケーションの場合、このサンプル コードでホストからデバイス メモリにデータを問題なく転送できますが、パフォーマンスおよび詳細な制御を最大にするために従う必要のあるコーディング プラクティスも多くあります。

clEnqueueMigrateMemObjects の使用

重要: エンベデッド プラットフォームの場合、CL_MEM_USE_HOST_PTRclEnqueueMigrateMemObjects を適用できません。エンベデッド プラットフォーム ユーザーは、clEnqueueMapBuffer の使用 で説明するように clEnqueueMapBuffer を使用する必要があります。

OpenCL フレームワークには、ホストとデバイス間でデータを転送するための API が多く含まれます。通常は、clEnqueueWriteBuffer および clEnqueueReadBuffer などのデータ移動 API でメモリ オブジェクトがキューに追加された後に暗示的にデバイスへ移動されますが、そのタイミングは決められないので、ホスト アプリケーションがメモリ オブジェクトの移動とデータで実行された計算を同期しにくくなります。

ザイリンクスでは、パフォーマンスを向上するために、clEnqueueWriteBuffer または clEnqueueReadBuffer の代わりに clEnqueueMigrateMemObjects を使用することをお勧めしています。メモリ移動が依存コマンドよりも前に明示的に実行されるようにできます。これにより、ホスト アプリケーションが通常のコマンド キューのスケジュールを使用して次の別のコマンドの準備ができるように、前もってメモリ オブジェクトの関連付けを変更できます。また、メモリ オブジェクトが必要となる前にその配置とその他の関連しない演算をオーバーラップさせて、発生する可能性のある転送レイテンシを隠したり、削減することもできます。clEnqueueMigrateMemObjects と関連するイベントが complete とマークされると、メモリ オブジェクトが問題なく移行できたことがホスト プログラムに伝わります。

cl_mem オブジェクトには、ホスト側ポインターとデバイス側ポインターの主に 2 つのポインターがあります。デバイス側ポインターは、カーネルが動作を開始する前に、デバイス側のメモリに暗示的に割り当てられる (たとえば、デバイスのグローバル メモリ内の特定位置に割り当てられる) ので、バッファーがデバイスに含まれるようになりますが、clEnqueueMigrateMemObjects を使用すると、この割り当てとデータ転送が前もって、カーネル実行よりもかなり前に発生するようになります。こうすることで、カーネルがまだ前のデータ セットを処理している間に次のトランザクションのデータ転送を実行して、連続するカーネル実行のデータ転送レイテンシを隠すことができるので、ホストが同じカーネルを何度も実行する場合に、「ソフトウェア パイプライン」をイネーブルにするのに特に役立ちます。

ヒント: また、clEnqueueMigrateMemObjects には、複数のメモリ オブジェクトを 1 つの API 呼び出しに移動できるという利点もあります。これにより、メモリ オブジェクトが複数ある場合のデータ転送のスケジューリングおよび関数呼び出しのオーバーヘッドが削減します。

次のコードは、clEnqueueMigrateMemObjects を使用したところを示しています。

int host_mem_ptr[MAX_LENGTH]; // host memory for input vector
      
// Fill the memory input
for(int i=0; i<MAX_LENGTH; i++) {
  host_mem_ptr[i] = <... >   
}

cl_mem dev_mem_ptr = clCreateBuffer(context,  
    				 CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
    				 sizeof(int) * number_of_words, host_mem_ptr, NULL); 

clSetKernelArg(kernel, 0, sizeof(cl_mem), &dev_mem_ptr); 

err = clEnqueueMigrateMemObjects(commands, 1, dev_mem_ptr, 0, 0, 
	  NULL, NULL);

ページ アライメントされたホスト メモリの割り当て

重要: このトピックは、エンベデッド プラットフォームには該当しません。エンベデッド プラットフォーム ユーザーは、clEnqueueMapBuffer の使用 で説明するように clEnqueueMapBuffer を使用する必要があります。

XRT は、4K 境界のメモリ空間を割り当てて、内部メモリを管理します。ホスト メモリのポインターが ページ境界に揃っていない場合、XRT が揃うように memcpy を追加で実行します。このため、ホスト メモリ ポインターを 4 K 境界に揃えて、メモリ コピー操作が余分に実行されないようにしてください。

次に、ホストのメモリ空間に malloc ではなく posix_memalign を使用する例を示します。

int *host_mem_ptr; // = (int*) malloc(MAX_LENGTH*sizeof(int));
// Aligning memory in 4K boundary
posix_memalign(&host_mem_ptr,4096,MAX_LENGTH*sizeof(int)); 
 
// Fill the memory input       
for(int i=0; i<MAX_LENGTH; i++) {
  host_mem_ptr[i] = <... >   
}

cl_mem dev_mem_ptr = clCreateBuffer(context, 
				     CL_MEM_READ_WRITE ,  
  			       sizeof(int) * number_of_words, host_mem_ptr, NULL); 

err = clEnqueueMigrateMemObjects(commands, 1, dev_mem_ptr, 0, 0, 
      NULL, NULL);

clEnqueueMapBuffer の使用

バッファーを作成して管理するには、clEnqueueMapBuffer を使用する方法もあります。この方法を使用する場合、4K 境界に揃えられたホスト空間ポインターを作成する必要はありません。clEnqueueMapBuffer API は指定したバッファーをマップして XRT で作成されたポインターをこのマップした領域に戻します。この後、ホスト側のポインターをユーザーのデータで埋めて、clEnqueueMigrateMemObject でデータをデバイスに送信したり受信したりします。次はその具体例です。


// Two cl_mem buffer, for read and write by kernel
cl_mem dev_mem_read_ptr = clCreateBuffer(context,  
    				 CL_MEM_READ_ONLY,
    				 sizeof(int) * number_of_words, NULL, NULL); 

cl_mem dev_mem_write_ptr = clCreateBuffer(context,  
    				 CL_MEM_WRITE_ONLY,
    				 sizeof(int) * number_of_words, NULL, NULL); 


// Setting arguments
clSetKernelArg(kernel, 0, sizeof(cl_mem), &dev_mem_read_ptr); 
clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_mem_write_ptr); 

// Get Host side pointer of the cl_mem buffer object
auto host_write_ptr = clEnqueueMapBuffer(queue,dev_mem_read_ptr,true,CL_MAP_WRITE,0,bytes,0,nullptr,nullptr,&err);
auto host_read_ptr = clEnqueueMapBuffer(queue,dev_mem_write_ptr,true,CL_MAP_READ,0,bytes,0,nullptr,nullptr,&err);

// Fill up the host_write_ptr to send the data to the FPGA

for(int i=0; i< MAX; i++) {
    host_write_ptr[i] = <.... > 
}

// Migrate
cl_mem mems[2] = {host_write_ptr,host_read_ptr};
clEnqueueMigrateMemObjects(queue,2,mems,0,0,nullptr,&migrate_event));

// Schedule the kernel
clEnqueueTask(queue,kernel,1,&migrate_event,&enqueue_event);

// Migrate data back to host
clEnqueueMigrateMemObjects(queue, 1, &dev_mem_write_ptr, 
                           CL_MIGRATE_MEM_OBJECT_HOST,1,&enqueue_event, &data_read_event);
     
clWaitForEvents(1,&data_read_event);

// Now use the data from the host_read_ptr

デバイスでのバッファーの割り当て

デフォルトでは、カーネルがリンクされる場合、すべてのカーネルからのメモリ インターフェイスが 1 つのデフォルトのグローバル メモリ バンクに接続されます。そのため、グローバル メモリ バンクとデータを転送できるのは一度に 1 つの計算ユニット (CU) のみになり、アプリケーションの全体的なパフォーマンスが制限されます。FPGA にグローバル メモリ バンクが 1 つしか含まれていない場合は、これしかオプションがありません。デバイスに複数のグローバル メモリ バンクが含まれる場合は、カーネル リンク中にそのグローバル メモリ バンクの接続をカスタマイズして、デフォルトの接続を変更できます。詳細は、カーネル ポートのグローバル メモリへのマップ を参照してください。全体的なパフォーマンスは、異なるカーネルまたは計算ユニットに対して別のメモリ バンクを使用して、複数のカーネル メモリ インターフェイスが同時にデータを読み出しおよび書き込みできるようにすると改善します。

重要: XRT はカーネルのメモリ接続を検出して、データをホスト プログラムからカーネルの正しいメモリ位置に送信する必要があります。XRT では、clSetKernelArgs がバッファーのエンキュー操作 (たとえば clEnqueueMigrateMemObject) よりも前に使用されると、自動的にカーネル バイナリ ファイルからバッファー位置を検出するようになっています。

サブバッファー

まれにですが、サブバッファーを使用すると、特殊な状況で役に立つことがあります。次のセクションでは、サブバッファーを使用すると利点がある例を示します。

デバイス バッファーから特定部分を読み込む方法

カーネルへの入力によって、出力されるデータ量が異なることがあります。たとえば、入力データ パターンによって、出力サイズが変更する圧縮エンジンがあります。ホストは clEnqueueMigrateMemObjects を使用して出力バッファー全体を読み出すことはできますが、必要以上のメモリ転送が発生してしまうので、最適な方法ではありません。理想的なのは、ホストがカーネルの書き込んだのと同じ量のデータだけを読み出す方法です。

たとえば、カーネルが出力データの書き込み開始時点で出力データ量を書き込むようにするという方法があります。ホスト アプリケーションは clEnqueueReadBuffer を 2 回使用できます (1 回目は返されたデータ量を読み出して、2 回目で最初の読み出しからの情報に基づいてカーネルから返されたのと同じ量のデータを読み出します)。
clEnqueueReadBuffer(command_queue,device_write_ptr, CL_FALSE, 0, sizeof(int) * 1, 
                    &kernel_write_size, 0, nullptr, &size_read_event);
clEnqueueReadBuffer(command_queue,device_write_ptr, CL_FALSE, DATA_READ_OFFSET, 
                    kernel_write_size, host_ptr, 1, &size_read_event, &data_read_event);
clEnqueueMigrateMemObject (clEnqueueReadBuffer または clEnqueueWriteBuffer よりも推奨) を使用すると、サブバッファーを使用して同様の方法を使用できます。次のコードはその具体例です。
ヒント: コード例は、この概念を示すことを目的としているので、コマンドの一部しか表示していません。
//Create a small sub-buffer to read the quantity of data
cl_buffer_region buffer_info_1={0,1*sizeof(int)}; 
cl_mem size_info = clCreateSubBuffer (device_write_ptr, CL_MEM_WRITE_ONLY, 
      CL_BUFFER_CREATE_TYPE_REGION, &buffer_info_1, &err);

// Map the sub-buffer into the host space
auto size_info_host_ptr = clEnqueueMapBuffer(queue, size_info,,,, );

// Read only the sub-buffer portion
clEnqueueMigrateMemObjects(queue, 1, &size_info, CL_MIGRATE_MEM_OBJECT_HOST,,,);
                          
// Retrive size information from the already mapped size_info_host_ptr
kernel_write_size = ........... 

// Create sub-buffer to read the required amount of data     
cl_buffer_region buffer_info_2={DATA_READ_OFFSET, kernel_write_size};
cl_mem  buffer_seg = clCreateSubBuffer (device_write_ptr, CL_MEM_WRITE_ONLY, 
      CL_BUFFER_CREATE_TYPE_REGION, &buffer_info_2,&err);

// Map the subbuffer into the host space
auto read_mem_host_ptr = clEnqueueMapBuffer(queue, buffer_seg,,,);

// Migrate the subbuffer
clEnqueueMigrateMemObjects(queue, 1, &buffer_seg, CL_MIGRATE_MEM_OBJECT_HOST,,,);

// Now use the read data from already mapped read_mem_host_ptr
複数のメモリ ポートまたは複数のカーネルで共有されるデバイス バッファー

カーネルのメモリ ポートが少量のデータしか必要としないことがありますが、小さなサイズのバッファーを管理し、少量のデータを転送すると、アプリケーションでパフォーマンスに問題が発生することもあります。このため、ホスト プログラムがより大きなサイズのバッファーを作成して、それを小さなサブバッファーに分割するという方法があります。カーネル引数の設定 に説明するように、サブバッファーはそれぞれ少量のデータを必要とするメモリ ポートごとにカーネル引数を割り当てます。

サブバッファーが作成されたら、標準的なバッファーと同様、ホスト コードで使用できるようになります。これにより、XRT が複数の小型バッファーで複数のトランザクションを実行するのではなく、大型のバッファーで 1 つのトランザクションを処理するようになるので、パフォーマンスが改善することがあります。

カーネル実行

ホスト アプリケーションで必要とされる演算負荷の高いタスクは、1 つのカーネル内で定義でき、そのカーネルを一度実行するだけでデータ範囲全体が処理されるようにできることがよくあります。複数のカーネル実行に関連するオーバーヘッドがある場合、1 つのモノリシック カーネルでパフォーマンスが改善されることがあります。この場合、カーネルを 1 回実行するだけでデータの範囲全体が処理されますが、並列処理 (したがってアクセラレーション) はカーネル ハードウェア内の FPGA で達成されます。カーネルは命令レベルの並列処理 (ループ パイプライン) および関数レベルの並列処理 (データフロー) などのさまざまな方法を使用して並列処理を達成します。これらのさまざまなカーネル コーディング手法については、C/C ++ カーネル を参照してください。

カーネルが FPGA の 1 つのハードウェア インスタンス (または CU) にコンパイルされる場合、次のように clEnqueueTask を使用するのがカーネルを実行する最も簡単な方法です。

err = clEnqueueTask(commands, kernel, 0, NULL, NULL);

XRT はワークロード (カーネル引数を使用し、OpenCL バッファーを介して渡されるデータ) をスケジュールし、さらに ザイリンクス FPGA 上のアクセラレータで実行されるようにカーネル タスクをスケジュールします。

重要: clEnqueueNDRangeKernel の使用はサポートされますが (OpenCL カーネルのみ)、ザイリンクスでは clEnqueueTask を使用することをお勧めしています。

ただし、カーネルを実行するのに 1 つの clEnqueueTask を使用するのは、さまざまな現実的な理由から、常に実現可能なわけではありません。たとえば、カーネル コードですべての演算負荷の高いタスクを 1 つの実行で処理されるようにすると、大きく複雑になりすぎて、最適化ができないことがあります。また、ホストがデータを受信するのに時間がかかったり、すべてのデータを一度に処理できないこともあります。このため、次のセクションに示すように、状況およびアプリケーションによって、データおよびタスクを複数の clEnqueueTask コマンドに分割する方法があります。

次のトピックでは、カーネルを実行したり、複数のカーネルを実行したり、アクセラレータ上の同じカーネルの複数インスタンスを実行するのに使用可能なさまざまな手法について説明します。

異なるカーネル使用したタスクの並列処理

ホスト アプリケーションの必要とする計算負荷の高いタスクは、複数の異なる カーネルに分割して、FPGA 上で異なるタスクを並列で実行するようにできることがあります。たとえば、複数の clEnqueueTask コマンドを順不同のコマンド キューで使用すると、異なるタスクを複数のカーネルで並列に実行できます。これにより、FPGA でタスクの並列処理ができるようになります。

空間的なデータ並列処理: 計算ユニット数を増加

ホスト アプリケーションの必要とする計算負荷の高いタスクでは、FPGA でデータ並列処理を達成するために、同じカーネルの複数のハードウェア インスタンスまたは計算ユニットでデータを処理できることがあります。1 つのカーネルが複数の CU にコンパイルされた場合、clEnqueueTask コマンドを順不同コマンドで複数回呼び出して、データ並列処理をイネーブルにします。clEnqueueTask を呼び出すたびに、異なる CU に含まれるデータのワークロードがスケジュールされて、並列処理されるようになります。

時間的なデータ並列処理: ホストからカーネルへのデータフロー

計算ユニットで処理されたデータがカーネルの 1 つの処理段階から次の処理段階へ渡されることがあります。この場合、カーネルの最初の段階では新しいデータ セットをいつでも自由に処理できます。つまり、工場の組み立てラインのように、カーネルが新しいデータを受信している間に元のデータがライン上を移動できます。

たとえば、カーネルに FPGA 上の 1 つのカーネルしか含まれない場合、ホスト アプリケーションが別のデータ セットを使用してカーネルを何度もエンキューするとします。clEnqueueMigrateMemObjects の使用 に示すように、ホスト アプリケーションはデータをカーネル実行前にデバイスのグローバル メモリに送信できるので、カーネル実行でデータ転送レイテンシを隠することで「ソフトウェア パイプライン」をイネーブルにできます。

ただし、デフォルトでは、カーネルは現在のデータ セットの処理を終了してからのみ、新しいデータ セットを処理し始めることができるようになっています。clEnqueueMigrateMemObject を使用すると、データ転送時間を隠すことはできますが、カーネル実行はシーケンシャルのままです。

ホストからカーネルへのデータフローをイネーブルにすると、カーネルがまだ前のデータ セットを処理している間に新しいデータ セットでカーネルを再開して、パフォーマンスをさらに改善できます。ホストからカーネルのデータフローのイネーブル で説明するように、カーネルは ap_ctrl_chain インターフェイスをインプリメントし、段階的にデータ処理を許可するように書き込まれる必要があります。この場合、XRT はカーネルが新しいデータを受信できるようになると即座に再開するので、複数のカーネル実行がオーバーラップします。ただし、カーネルが新しいデータの受信準備完了後すぐに再開できるように、ホスト プログラムは要求でいっぱいになったコマンド キューを維持する必要があります。

次の図は、ホストからカーネルへのデータフローの概念を示しています。

1: ホストからカーネルのデータフロー

カーネルで開始から終了までデータ セットを処理するのに時間がかかるほど、ホストからカーネルへのデータフローが使用される機会が増えて、パフォーマンスが改善します。カーネルが 1 つのデータ セットを終了するまで待つよりも、単にカーネルが次のデータ セットを処理し始める準備ができるのを待ちます。これにより、「時間的な並列処理」ができるようになり、同じカーネルの異なる段階で複数の clEnqueueTask コマンドからの異なるデータ セットをパイプライン処理できるようになります。

高度なデザインの場合は、データを処理するのに複数の CU を使用する空間的な並列処理とホストからカーネルへのデータフローを使用した時間的な並列処理を効率的に組み合わせて使用し、各計算ユニットのカーネル実行をオーバーラップさせます。

重要: エンベデッド プロセッサ プラットフォームでは、ホストからカーネルへのデータフロー機能はサポートされません。

対称および非対称の計算ユニット

複数のカーネル インスタンスの作成 で説明したように、カーネル リンク プロセスの際には、1 つのカーネルの複数の計算ユニット (CU) を FPGA にインスタンシエートできます。CU は同じカーネルのほかの CU に対して対称または非対称になります。

対称
CU は同じ connectivity.sp オプションが使用されると対称となるので、それらの CU のグローバル メモリへの接続はまったく同じになります。この結果、ザイリンクス ランタイム (XRT) がそれらを交互に使用できるようになります。clEnqueueTask を呼び出すと、対称 CU のグループのいずれかのインスタンスが開始されます。
非対称
CU に同じ connectivity.sp オプションが使用されていないと、それらは非対称となるので、グローバル メモリへの接続は同じにはなりません。同じ入力および出力 バッファー設定を使用した場合、XRT で非同期 CU を交互に実行できません。
カーネル ハンドルおよび計算ユニット

clSetKernelArg が指定したカーネル オブジェクトに対して最初に呼び出されると、XRT ではこのカーネルの後に続く実行に対称 CU のグループが識別されます。そのカーネル用に clEnqueueTask が呼び出されると、そのグループの対称 CU をどれでも使用してタスクを処理できるようになります。

指定したカーネルの CU すべてが対称の場合は、1 つのカーネル オブジェクトだけでこれらのどの CU にでもアクセスできますが、非対称 CU がある場合は、ホスト アプリケーションが非同期 CU の各グループに対して一意のカーネル オブジェクトを作成する必要があります。この場合、clEnqueueTask への呼び出しでタスクに使用するカーネル オブジェクトを指定する必要があり、そのカーネルに一致する CU が XRT で使用できるようになります。

特定の計算ユニット用のカーネル オブジェクトの作成

特定の計算ユニットに関連するカーネルを作成する場合、 clCreateKernel コマンドを使用すると、カーネル オブジェクトがホスト プログラムで作成されるときに CU を指定できます。このコマンドの構文は、次のとおりです。

// Create kernel object only for a specific compute unit 
cl_kernel kernelA = clCreateKernel(program,"<kernel_name>:{compute_unit_name}",&err);
// Create a kernel object for two specific compute units 
cl_kernel kernelB = clCreateKernel(program, "<kernel_name>:{CU1,CU2}", &err);
重要: 複数のカーネル インスタンスの作成 で説明するように、CU 数はリンク中に v++ コマンドで使用されるコンフィギュレーション ファイルの connectivity.nk オプションで指定します。このため、カーネル オブジェクトを作成またはエンキューするには、ホスト プログラムで何を指定したとしても、リンク中に使用されるコンフィギュレーション ファイルで指定したオプションと同じにする必要があります。

この場合、カーネルが作成される際、ザイリンクス ランタイムで指定した CU または CU のグループのカーネル ハンドル (kernelAkernelB) が識別されます。これで、ホスト プログラム内の clEnqueueTask を使用する際に、どのカーネル コンフィギュレーションまたは特定の CU インスタンスを使用するか制御できるようになります。これは、非対称 CU がある場合や、CU のロードおよび優先度管理を実行する場合に役立つことがあります。

計算ユニット名を使用して非対称計算ユニットすべてのハンドルを取得

カーネルが複数の対称ではない CU をインスタンシエートすると、clCreateKernel コマンドに CU 名を指定して、異なる CU グループを作成します。この場合、ホスト プログラムは clCreateKernel で返される cl_kernel ハンドルを使用して特定の CU グループを参照できます。

次の例では、mykernel カーネルに 5 つの CU (K1、K2、K3、K4、K5) が含まれます。K1、K2、K3 計算ユニットは、デバイス上に対称接続を持つ 1 つの対称グループです。同様に、K4 および K5 計算ユニットは 2 つ目の対称 CU のグループです。次のコード セグメントは、cl_kernel ハンドルを使用して特定の CU グループを指定するところを示しています。

// Kernel handle for Symmetrical compute unit group 1: K1,K2,K3
cl_kernel kernelA = clCreateKernel(program,"mykernel:{K1,K2,K3}",&err);

for(i=0; i<3; i++) {
  // Creating buffers for the kernel_handle1
  .....
  // Setting kernel arguments for kernel_handle1
  .....
  // Enqueue buffers for the kernel_handle1
  .....
  // Possible candidates of the executions K1,K2 or K3
  clEnqueueTask(commands, kernelA, 0, NULL, NULL); 
  //
}

// Kernel handle for Symmetrical compute unit group 1: K4, K5
cl_kernel kernelB = clCreateKernel(program,"mykernel:{K4,K5}",&err);

for(int i=0; i<2; i++) {
  // Creating buffers for the kernel_handle2
  .....
  // Setting kernel arguments for kernel_handle2
  .....
  // Enqueue buffers for the kernel_handle2
  .....
  // Possible candidates of the executions K4 or K5
  clEnqueueTask(commands, kernelB, 0, NULL, NULL);
}

イベントの同期化

OpenCL エンキュー ベースの API 呼び出しはすべて非同期です。これらのコマンドはコマンド キューに追加されるとすぐに返されます。ホスト プログラムが結果を待ったり、コマンド間の依存を解決する間に停止するようにするには、clFinish または clWaitForEvents などの API 呼び出しを使用すると、ホスト プログラムの実行をブロックできます。

次は、clFinish および clWaitForEvents のコード例です。

err = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);
// Execution will wait here until all commands in the command queue are finished
clFinish(command_queue); 

// Create event, read memory from device, wait for read to complete, verify results
cl_event readevent;
// host memory for output vector
int host_mem_output_ptr[MAX_LENGTH]; 
//Enqueue ReadBuffer, with associated event object 
clEnqueueReadBuffer(command_queue, dev_mem_ptr, CL_TRUE, 0, sizeof(int) * number_of_words, 
  host_mem_output_ptr, 0, NULL, &readevent );
// Wait for clEnqueueReadBuffer event to finish
clWaitForEvents(1, &readevent); 
// After read is complete, verify results
... 

上記の例でコマンドがどのように使用されているかに注目してください。

  1. clFinish API は、カーネルの実行が終了するまでホストが実行されないようにするために使用されています。このようにしないと、ホストが FPGA バッファーからデータを読み出すのが早すぎ、無効なデータが読み出される可能性があります。
  2. FPGA メモリからローカル ホスト マシンへのデータ転送は、clEnqueueReadBuffer を使用して実行されます。clEnqueueReadBuffer の最後の引数は、この特定の読み出しコマンドを識別して、イベントをクエリするのに使用可能な (またはこの特定コマンドが終了するまで待機する) イベント オブジェクトを返します。clWaitForEvents コマンドはその 1 つのイベント (readevent) を指定して、データ転送が終了するのを待ってから、データを検証します。

後処理および FPGA のクリーンアップ

ホスト コードの最後には、適切な解放関数を使用して、割り当てられたリソースをすべて解放する必要があります。リソースが正しく開放されない場合は、Vitis コア開発キットで正しいパフォーマンスに関連したプロファイルおよび解析レポートを生成できないことがあります。
clReleaseCommandQueue(Command_Queue);
clReleaseContext(Context);
clReleaseDevice(Target_Device_ID); 
clReleaseKernel(Kernel);
clReleaseProgram(Program);
free(Platform_IDs);
free(Device_IDs);

まとめ

前のトピックで説明したように、Vitis コア開発キットで推奨されるホスト プログラムのコーディング スタイルは次のとおりです。

  1. 必要に応じて、デバッグ用に各 OpenCL API 呼び出しの後にエラー チェックを追加します。
  2. Vitis コア開発キットでは、1 つまたは複数のカーネルは個別にコンパイル/リンクされて、.xclbin ファイルがビルドされます。カーネル バイナリから cl_program オブジェクトをビルドするには、clCreateProgramWithBinary API を使用します。
  3. バッファーのエンキュー操作前にカーネル引数 (clSetKernelArg) を設定するバッファーを使用します。
  4. ホスト コードと FPGA の間のデータ転送には clEnqueueMigrateMemObjects または clEnqueueMapBuffer を使用します。
  5. posix_memalign を使用してホスト メモリ ポインターを 4K 境界に合わせます (PCIe ベース プラットフォームに使用可能)。
  6. FPGA で同時実行コマンドを実行する順不同コマンド キューを使用することをお勧めします。
  7. clEnqueueNDRangeKernel を使用してワークロードを分割するのではなく、clEnqueueTask を使用してワークロード全体を実行します。
  8. イベント同期化コマンドの clFinish および clWaitForEvents を使用して、非同期 OpenCL API 呼び出しの依存を解決します。
  9. 終了したらすべての OpenCL 割り当てのリソースを解放します。