Programming the Host Application

Sub-devices

In the SDAccel environment, sometimes devices contain multiple kernel instances, of a single kernel or of different kernels. The OpenCL API clCreateSubDevices allows the host code to divide the device into multiple sub-devices containing one kernel instance per sub-device. Currently SDAccel environment supports equally divided sub-devices each containing only one kernel instance.

  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 
  });

Currently, if a kernel has multiple hardware instances (can be specified during the kernel compilation phase), the SDAccel environment execution model assumes all those hardware instances have the same global memory connectivity. If not, then you need to use sub-devices to allocate separate cl_kernel for each of those hardware instances.

One or more command queues for each device is created using the clCreateCommandQueue API. The FPGA device can contain multiple kernels. When developing the host application, there are two main programming approaches to execute kernels on a device:

1. Single out-of-order command queue: Multiple kernel executions can be requested through the same command queue. The SDAccel runtime environment dispatches those kernels as soon as possible in any order allowing concurrent kernel execution on the FPGA.
2. Multiple in-order command queue: Each kernel execution will be requested from different in-order command queues. In such cases, the SDAccel runtime environment can dispatch kernels from any command queue with the intention of improving performance by running them concurrently on the FPGA.

// 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);

Interactions between the host application and kernels rely on transferring data to and from global memory in the device. The simplest way to send data back and forth from the FPGA is using clCreateBuffer, clEnqueueWriteBuffer and clEnqueueReadBuffer commands. The following code example demonstrates this:

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);

For the majority of applications the example code above would be sufficient to transfer data from the host to the device memory. However, there are a number of coding practices you should adopt in order to maximize performance and fine-grain control.

Using clEnqueueMigrateMemObjects

Another consideration when transferring data is using clEnqueueMigrateMemObjects instead of clEnqueueWriteBuffer or clEnqueueReadBuffer to improve the performance. Typically, memory objects are implicitly migrated to a device for enqueued kernels. Using this API call results in data transfer ahead of kernel execution to reduce latency, particularly when a kernel is called multiple times.

The following code example is modified to use 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_ext_ptr_t d_bank0_ext;
  d_bank0_ext.flags = XCL_MEM_DDR_BANK0;
  d_bank0_ext.obj = host_mem_ptr; 
  d_bank0_ext.param = 0;

cl_mem dev_mem_ptr = clCreateBuffer(context,  
    				 CL_MEM_READ_WRITE | CL_MEM_EXT_PTR_XILINX,
    				 sizeof(int) * number_of_words, &d_bank0_ext, NULL); 

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

Using posix_memalign for Host Memory Space

SDAccel runtime allocates the memory space in 4K boundary for internal memory management. If the host memory pointer is not aligned to a 4K word boundary, the runtime performs extra memcpy to make it aligned. It does not significantly impact performance, but you should align the host memory pointer with the 4K boundary to follow the SDAccel runtime memory management.

The following is an example of how posix_memalign is used instead of malloc for the host memory space pointer.

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_ext_ptr_t d_bank0_ext;
  d_bank0_ext.flags = XCL_MEM_DDR_BANK0;
  d_bank0_ext.obj = host_mem_ptr; 
  d_bank0_ext.param = 0;

cl_mem dev_mem_ptr = clCreateBuffer(context, 
				     CL_MEM_READ_WRITE | CL_MEM_EXT_PTR_XILINX,  
  			       sizeof(int) * number_of_words, &d_bank0_ext, NULL); 

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

Enhanced Buffer Allocation

By default, all the memory interfaces from all the kernels are connected to a single global memory bank when kernels are linked. As a result, only one memory interface can transfer data to and from the global memory bank at a time, limiting the overall performance of the application. If the FPGA device contains only one global memory bank, this is the only option. However, if the device contains multiple global memory banks, you can customize the global memory bank connections by modifying the default connection. This topic is discussed in greater detail in Customization of DDR Bank to Kernel Connection. This improves overall performance by enabling multiple kernel memory interfaces to concurrently read and write data from separate global memory banks.

When kernel ports are mapped to memory banks other than the default one, it is necessary to use the enhanced buffer allocation pattern when creating the OpenCL buffers.

The enhanced buffer allocation pattern uses a Xilinx vendor extension, cl_mem_ext_ptr_t, pointer to help the Xilinx runtime determine which global memory bank the buffer should be allocated.

The cl_mem_ext_ptr_t type is a struct as defined below:

typedef struct{
     unsigned flags;
     void *obj;
     void *param;
  } 
  cl_mem_ext_ptr_t;

Use the explicit bank name method to operate cl_mem_ext_ptr_t for enhanced buffer allocation.

Explicit Bank Name Method

In this approach, the struct field flags is used to denote the DDR bank (XCL_MEM_DDR_BANK1, XCL_MEM_DDR_BANK2, etc.). The struct field param should not be used and set to NULL.

The following code example uses cl_mem_ext_ptr_t to assign the device buffer to DDR Bank 2.

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_ext_ptr_t d_bank0_ext;
  d_bank0_ext.flags = XCL_MEM_DDR_BANK2;
  d_bank0_ext.obj = NULL; 
  d_bank0_ext.param = 0;

cl_mem dev_mem_ptr = clCreateBuffer(context,
				   CL_MEM_READ_WRITE | CL_MEM_EXT_PTR_XILINX,
   			     sizeof(int) * number_of_words, &d_bank0_ext, NULL); 

err = clEnqueueWriteBuffer(commands, dev_mem_ptr, CL_TRUE, 0,
      sizeof(int) * number_of_words, host_mem_ptr, 0, NULL, NULL);

var_ext.flags =  <some integer> | XCL_MEM_TOPOLOGY

Identifying the kernels

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

Setting Kernel Arguments

int err = 0;
// Setting up scalar arguments
cl_uint scalar_arg_image_width = 3840;
err |= clSetKernelArg(kernel, 0, sizeof(cl_uint), &scaler_arg_image_width); 
cl_uint scaler_arg_image_height = 2160; 
err |= clSetKernelArg(kernel, 1, sizeof(cl_uint), &scaler_arg_image_height); 
    
// Setting up buffer arguments
err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &dev_mem_ptr0);
err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &dev_mem_ptr1);
err |= clSetKernelArg(kernel, 4, sizeof(cl_mem), &dev_mem_ptr2);

Enqueing the Kernels

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

Just like all the enqueue commands, the clEnqueueTask and clEnqueueNDRangeKernel are asynchronous in nature. The host code continues executing without waiting for the kernel computation to complete on the FPGA device. This allows the host program to execute more kernels, either the same kernel multiple times over a different set of data, or different kernel. After finishing its work, the kernel writes the result data to the global memory bank. This data is read back to the host memory space by using clEnqueueReadBuffer or the clEnqueueMigrateMemObjects command.