xilinx.md

# Xilinx Accelerator Platform

The first step to use Xilinx accelerators is to initialize Vitis (compiler) and XRT (runtime) environments.

```console
$ . /tools/Xilinx/Vitis/2023.1/settings64.sh
$ . /opt/xilinx/xrt/setup.sh
```

## Platform Level Accelerator Management

This should allow to examine current platform using `xbutil examine`,
which should output user-level information about XRT platform and list available devices

```
$ xbutil examine
System Configuration
  OS Name              : Linux
  Release              : 4.18.0-477.27.1.el8_8.x86_64
  Version              : #1 SMP Thu Aug 31 10:29:22 EDT 2023
  Machine              : x86_64
  CPU Cores            : 64
  Memory               : 257145 MB
  Distribution         : Red Hat Enterprise Linux 8.8 (Ootpa)
  GLIBC                : 2.28
  Model                : ProLiant XL675d Gen10 Plus

XRT
  Version              : 2.16.0
  Branch               : master
  Hash                 : f2524a2fcbbabd969db19abf4d835c24379e390d
  Hash Date            : 2023-10-11 14:01:19
  XOCL                 : 2.16.0, f2524a2fcbbabd969db19abf4d835c24379e390d
  XCLMGMT              : 2.16.0, f2524a2fcbbabd969db19abf4d835c24379e390d

Devices present
BDF             :  Shell                            Logic UUID                            Device ID       Device Ready*
-------------------------------------------------------------------------------------------------------------------------
[0000:88:00.1]  :  xilinx_u280_gen3x16_xdma_base_1  283BAB8F-654D-8674-968F-4DA57F7FA5D7  user(inst=132)  Yes
[0000:8c:00.1]  :  xilinx_u280_gen3x16_xdma_base_1  283BAB8F-654D-8674-968F-4DA57F7FA5D7  user(inst=133)  Yes


* Devices that are not ready will have reduced functionality when using XRT tools
```

Here two Xilinx Alveo u280 accelerators (`0000:88:00.1` and `0000:8c:00.1`) are available.
The `xbutil` can be also used to query additional information about specific device using its BDF address

```console
$ xbutil examine -d "0000:88:00.1"

-------------------------------------------------
[0000:88:00.1] : xilinx_u280_gen3x16_xdma_base_1
-------------------------------------------------
Platform
  XSA Name               : xilinx_u280_gen3x16_xdma_base_1
  Logic UUID             : 283BAB8F-654D-8674-968F-4DA57F7FA5D7
  FPGA Name              :
  JTAG ID Code           : 0x14b7d093
  DDR Size               : 0 Bytes
  DDR Count              : 0
  Mig Calibrated         : true
  P2P Status             : disabled
  Performance Mode       : not supported
  P2P IO space required  : 64 GB

Clocks
  DATA_CLK (Data)        : 300 MHz
  KERNEL_CLK (Kernel)    : 500 MHz
  hbm_aclk (System)      : 450 MHz

Mac Addresses            : 00:0A:35:0E:20:B0
                         : 00:0A:35:0E:20:B1

  Device Status: HEALTHY
  Hardware Context ID: 0
    Xclbin UUID: 6306D6AE-1D66-AEA7-B15D-446D4ECC53BD
    PL Compute Units
      Index  Name         Base Address  Usage  Status
      -------------------------------------------------
      0      vadd:vadd_1  0x800000      1      (IDLE)
```

Basic functionality of the device can be checked using `xbutil validate -d <BDF>` as

```console
$ xbutil validate -d "0000:88:00.1"
Validate Device           : [0000:88:00.1]
    Platform              : xilinx_u280_gen3x16_xdma_base_1
    SC Version            : 4.3.27
    Platform ID           : 283BAB8F-654D-8674-968F-4DA57F7FA5D7
-------------------------------------------------------------------------------
Test 1 [0000:88:00.1]     : aux-connection
    Test Status           : [PASSED]
-------------------------------------------------------------------------------
Test 2 [0000:88:00.1]     : pcie-link
    Test Status           : [PASSED]
-------------------------------------------------------------------------------
Test 3 [0000:88:00.1]     : sc-version
    Test Status           : [PASSED]
-------------------------------------------------------------------------------
Test 4 [0000:88:00.1]     : verify
    Test Status           : [PASSED]
-------------------------------------------------------------------------------
Test 5 [0000:88:00.1]     : dma
    Details               : Buffer size - '16 MB' Memory Tag - 'HBM[0]'
                            Host -> PCIe -> FPGA write bandwidth = 11988.9 MB/s
                            Host <- PCIe <- FPGA read bandwidth = 12571.2 MB/s
                            ...
    Test Status           : [PASSED]
-------------------------------------------------------------------------------
Test 6 [0000:88:00.1]     : iops
    Details               : IOPS: 387240(verify)
    Test Status           : [PASSED]
-------------------------------------------------------------------------------
Test 7 [0000:88:00.1]     : mem-bw
    Details               : Throughput (Type: DDR) (Bank count: 2) : 33932.9MB/s
                            Throughput of Memory Tag: DDR[0] is 16974.1MB/s
                            Throughput of Memory Tag: DDR[1] is 16974.2MB/s
                            Throughput (Type: HBM) (Bank count: 1) : 12383.7MB/s
    Test Status           : [PASSED]
-------------------------------------------------------------------------------
Test 8 [0000:88:00.1]     : p2p
Test 9 [0000:88:00.1]     : vcu
Test 10 [0000:88:00.1]    : aie
Test 11 [0000:88:00.1]    : ps-aie
Test 12 [0000:88:00.1]    : ps-pl-verify
Test 13 [0000:88:00.1]    : ps-verify
Test 14 [0000:88:00.1]    : ps-iops
```

Finally, the device can be reinitialized using `xbutil reset -d <BDF>` as

```console
$ xbutil reset -d "0000:88:00.1"
Performing 'HOT Reset' on '0000:88:00.1'
Are you sure you wish to proceed? [Y/n]: Y
Successfully reset Device[0000:88:00.1]
```

This can be useful to recover the device from states such as `HANGING`, reported by `xbutil examine -d <BDF>`.

## OpenCL Platform Level

The `clinfo` utility can be used to verify that the accelerator is visible to OpenCL

```console
$ clinfo
Number of platforms:                             2
  Platform Profile:                              FULL_PROFILE
  Platform Version:                              OpenCL 2.1 AMD-APP (3590.0)
  Platform Name:                                 AMD Accelerated Parallel Processing
  Platform Vendor:                               Advanced Micro Devices, Inc.
  Platform Extensions:                           cl_khr_icd cl_amd_event_callback
  Platform Profile:                              EMBEDDED_PROFILE
  Platform Version:                              OpenCL 1.0
  Platform Name:                                 Xilinx
  Platform Vendor:                               Xilinx
  Platform Extensions:                           cl_khr_icd
<...>
  Platform Name:                                 Xilinx
Number of devices:                               2
  Device Type:                                   CL_DEVICE_TYPE_ACCRLERATOR
  Vendor ID:                                     0h
  Max compute units:                             0
  Max work items dimensions:                     3
    Max work items[0]:                           4294967295
    Max work items[1]:                           4294967295
    Max work items[2]:                           4294967295
  Max work group size:                           4294967295
  Preferred vector width char:                   1
  Preferred vector width short:                  1
  Preferred vector width int:                    1
  Preferred vector width long:                   1
  Preferred vector width float:                  1
  Preferred vector width double:                 0
  Max clock frequency:                           0Mhz
  Address bits:                                  64
  Max memory allocation:                         4294967296
  Image support:                                 Yes
  Max number of images read arguments:           128
  Max number of images write arguments:          8
  Max image 2D width:                            8192
  Max image 2D height:                           8192
  Max image 3D width:                            2048
  Max image 3D height:                           2048
  Max image 3D depth:                            2048
  Max samplers within kernel:                    0
  Max size of kernel argument:                   2048
  Alignment (bits) of base address:              32768
  Minimum alignment (bytes) for any datatype:    128
  Single precision floating point capability
    Denorms:                                     No
    Quiet NaNs:                                  Yes
    Round to nearest even:                       Yes
    Round to zero:                               No
    Round to +ve and infinity:                   No
    IEEE754-2008 fused multiply-add:             No
  Cache type:                                    None
  Cache line size:                               64
  Cache size:                                    0
  Global memory size:                            0
  Constant buffer size:                          4194304
  Max number of constant args:                   8
  Local memory type:                             Scratchpad
  Local memory size:                             16384
  Error correction support:                      1
  Profiling timer resolution:                    1
  Device endianess:                              Little
  Available:                                     No
  Compiler available:                            No
  Execution capabilities:
    Execute OpenCL kernels:                      Yes
    Execute native function:                     No
  Queue on Host properties:
    Out-of-Order:                                Yes
    Profiling:                                   Yes
  Platform ID:                                   0x16fbae8
  Name:                                          xilinx_u280_gen3x16_xdma_base_1
  Vendor:                                        Xilinx
  Driver version:                                1.0
  Profile:                                       EMBEDDED_PROFILE
  Version:                                       OpenCL 1.0
<...>
```

which shows that both `Xilinx` platform and accelerator devices are present.

## Building Applications

To simplify the build process we define two environment variables `IT4I_PLATFORM` and `IT4I_BUILD_MODE`.
The first `IT4I_PLATFORM` denotes specific accelerator hardware such as `Alveo u250` or `Alveo u280`
and its configuration stored in (`*.xpfm` files).
The list of available platforms can be obtained using `platforminfo` utility:

```console
$ platforminfo -l
{
    "platforms": [
        {
            "baseName": "xilinx_u280_gen3x16_xdma_1_202211_1",
            "version": "202211.1",
            "type": "sdaccel",
            "dataCenter": "true",
            "embedded": "false",
            "externalHost": "true",
            "serverManaged": "true",
            "platformState": "impl",
            "usesPR": "true",
            "platformFile": "\/opt\/xilinx\/platforms\/xilinx_u280_gen3x16_xdma_1_202211_1\/xilinx_u280_gen3x16_xdma_1_202211_1.xpfm"
        },
        {
            "baseName": "xilinx_u250_gen3x16_xdma_4_1_202210_1",
            "version": "202210.1",
            "type": "sdaccel",
            "dataCenter": "true",
            "embedded": "false",
            "externalHost": "true",
            "serverManaged": "true",
            "platformState": "impl",
            "usesPR": "true",
            "platformFile": "\/opt\/xilinx\/platforms\/xilinx_u250_gen3x16_xdma_4_1_202210_1\/xilinx_u250_gen3x16_xdma_4_1_202210_1.xpfm"
        }
    ]
}
```

Here, `baseName` and potentially `platformFile` are of interest and either can be specified as value of `IT4I_PLATFORM`.
In this case we have platform files `xilinx_u280_gen3x16_xdma_1_202211_1` (Alveo u280) and `xilinx_u250_gen3x16_xdma_4_1_202210_1` (Alveo u250).

The `IT4I_BUILD_MODE` is used to specify build type (`hw`, `hw_emu` and `sw_emu`):

- `hw` performs full synthesis for the accelerator
- `hw_emu` allows to run both synthesis and emulation for debugging
- `sw_emu` compiles kernels only for emulation (doesn't require accelerator and allows much faster build)

For example to configure build for `Alveo u280` we set:

```console
$ export IT4I_PLATFORM=xilinx_u280_gen3x16_xdma_1_202211_1
```

### Software Emulation Mode

The software emulation mode is preferable for development as HLS synthesis is very time consuming. To build following applications in this mode we set:

```console
$ export IT4I_BUILD_MODE=sw_emu
```

and run each application with `XCL_EMULATION_MODE` set to `sw_emu`:

```
$ XCL_EMULATION_MODE=sw_emu <application>
```

### Hardware Synthesis Mode

!!! note
    The HLS of these simple applications **can take up to 2 hours** to finish.

To allow the application to utilize real hardware we have to synthetize FPGA design for the accelerator. This can be done by repeating same steps used to build kernels in emulation mode, but with `IT4I_BUILD_MODE` set to `hw` like so:

```console
$ export IT4I_BUILD_MODE=hw
```

the host application binary can be reused, but it has to be run without `XCL_EMULATION_MODE`:

```console
$ <application>
```

## Sample Applications

The first two samples illustrate two main approaches to building FPGA accelerated applications using Xilinx platform - **XRT** and **OpenCL**.
The final example combines **HIP** with **XRT** to show basics necessary to build application, which utilizes both GPU and FPGA accelerators.

### Using HLS and XRT

The applications are typically separated into host and accelerator/kernel side.
The following host-side code should be saved as `host.cpp`

```c++
/*
# Copyright (C) 2023, Advanced Micro Devices, Inc. All rights reserved.
# SPDX-License-Identifier: X11
*/
#include <iostream>
#include <cstring>

// XRT includes
#include "xrt/xrt_bo.h"
#include <experimental/xrt_xclbin.h>
#include "xrt/xrt_device.h"
#include "xrt/xrt_kernel.h"

#define DATA_SIZE 4096

int main(int argc, char** argv)
{
    if(argc != 2)
    {
        std::cout << "Usage: " << argv[0] << " <XCLBIN File>" << std::endl;
        return EXIT_FAILURE;
    }

    // Read settings
    std::string binaryFile = argv[1];
    int device_index = 0;

    std::cout << "Open the device" << device_index << std::endl;
    auto device = xrt::device(device_index);
    std::cout << "Load the xclbin " << binaryFile << std::endl;
    auto uuid = device.load_xclbin("./vadd.xclbin");

    size_t vector_size_bytes = sizeof(int) * DATA_SIZE;

    //auto krnl = xrt::kernel(device, uuid, "vadd");
    auto krnl = xrt::kernel(device, uuid, "vadd", xrt::kernel::cu_access_mode::exclusive);

    std::cout << "Allocate Buffer in Global Memory\n";
    auto boIn1 = xrt::bo(device, vector_size_bytes, krnl.group_id(0)); //Match kernel arguments to RTL kernel
    auto boIn2 = xrt::bo(device, vector_size_bytes, krnl.group_id(1));
    auto boOut = xrt::bo(device, vector_size_bytes, krnl.group_id(2));

    // Map the contents of the buffer object into host memory
    auto bo0_map = boIn1.map<int*>();
    auto bo1_map = boIn2.map<int*>();
    auto bo2_map = boOut.map<int*>();
    std::fill(bo0_map, bo0_map + DATA_SIZE, 0);
    std::fill(bo1_map, bo1_map + DATA_SIZE, 0);
    std::fill(bo2_map, bo2_map + DATA_SIZE, 0);

    // Create the test data
    int bufReference[DATA_SIZE];
    for (int i = 0; i < DATA_SIZE; ++i)
    {
        bo0_map[i] = i;
        bo1_map[i] = i;
        bufReference[i] = bo0_map[i] + bo1_map[i]; //Generate check data for validation
    }

    // Synchronize buffer content with device side
    std::cout << "synchronize input buffer data to device global memory\n";
    boIn1.sync(XCL_BO_SYNC_BO_TO_DEVICE);
    boIn2.sync(XCL_BO_SYNC_BO_TO_DEVICE);

    std::cout << "Execution of the kernel\n";
    auto run = krnl(boIn1, boIn2, boOut, DATA_SIZE); //DATA_SIZE=size
    run.wait();

    // Get the output;
    std::cout << "Get the output data from the device" << std::endl;
    boOut.sync(XCL_BO_SYNC_BO_FROM_DEVICE);

    // Validate results
    if (std::memcmp(bo2_map, bufReference, vector_size_bytes))
        throw std::runtime_error("Value read back does not match reference");

    std::cout << "TEST PASSED\n";
    return 0;
}
```

The host-side code can now be compiled using GCC toolchain as:

```console
$ g++ host.cpp -I$XILINX_XRT/include -I$XILINX_VIVADO/include -L$XILINX_XRT/lib -lxrt_coreutil -o host
```

The accelerator side (simple vector-add kernel) should be saved as `vadd.cpp`.

```c++
/*
# Copyright (C) 2023, Advanced Micro Devices, Inc. All rights reserved.
# SPDX-License-Identifier: X11
*/

extern "C" {
	void vadd(
	        const unsigned int *in1, // Read-Only Vector 1
	        const unsigned int *in2, // Read-Only Vector 2
	        unsigned int *out,       // Output Result
	        int size                 // Size in integer
	        )
	{
#pragma HLS INTERFACE m_axi port=in1 bundle=aximm1
#pragma HLS INTERFACE m_axi port=in2 bundle=aximm2
#pragma HLS INTERFACE m_axi port=out bundle=aximm1

	    for(int i = 0; i < size; ++i)
	    {
	        out[i] = in1[i] + in2[i];
	    }
	}
}
```

The accelerator-side code is build using Vitis `v++`.
This is two-step process, which either builds emulation binary or performs full HLS (depending on the value of `-t` argument).
The platform (specific accelerator) has to be also specified at this step (both for emulation and full HLS).

```console
$ v++ -c -t $IT4I_BUILD_MODE --platform $IT4I_PLATFORM -k vadd vadd.cpp -o vadd.xo
$ v++ -l -t $IT4I_BUILD_MODE --platform $IT4I_PLATFORM vadd.xo -o vadd.xclbin
```

This process should result in `vadd.xclbin`, which can be loaded by host-side application.

### Running the Application

With both host application and kernel binary at hand the application (in emulation mode) can be launched as

```console
$ XCL_EMULATION_MODE=sw_emu ./host vadd.xclbin
```

or with real hardware (having compiled kernels with `IT4I_BUILD_MODE=hw`)

```console
./host vadd.xclbin
```

## Using HLS and OpenCL

The host-side application code should be saved as `host.cpp`.
This application attempts to find `Xilinx` OpenCL platform in the system and selects first device in that platform.
The device is then configured with provided kernel binary.
Other than that the only difference to typical vector-add in OpenCL is use of `enqueueTask(...)` to launch the kernel
(compared to typical `enqueueNDRangeKernel`).

```c++
#include <iostream>
#include <fstream>
#include <iterator>
#include <vector>

#define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_MINIMUM_OPENCL_VERSION 120
#define CL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY 1
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS

#include <CL/cl2.hpp>
#include <CL/cl_ext_xilinx.h>

std::vector<unsigned char> read_binary_file(const std::string &filename)
{
    std::cout << "INFO: Reading " << filename << std::endl;
    std::ifstream file(filename, std::ios::binary);
    file.unsetf(std::ios::skipws);

    std::streampos file_size;
    file.seekg(0, std::ios::end);
    file_size = file.tellg();
    file.seekg(0, std::ios::beg);

    std::vector<unsigned char> data;
    data.reserve(file_size);
    data.insert(data.begin(),
        std::istream_iterator<unsigned char>(file),
        std::istream_iterator<unsigned char>());

    return data;
}

cl::Device select_device()
{
    std::vector<cl::Platform> platforms;
    cl::Platform::get(&platforms);
    cl::Platform platform;

    for(cl::Platform &p: platforms)
    {
        const std::string name = p.getInfo<CL_PLATFORM_NAME>();
        std::cout << "PLATFORM: " << name << std::endl;
        if(name == "Xilinx")
        {
            platform = p;
            break;
        }
    }

    if(platform == cl::Platform())
    {
        std::cout << "Xilinx platform not found!" << std::endl;
        exit(EXIT_FAILURE);
    }

    std::vector<cl::Device> devices;
    platform.getDevices(CL_DEVICE_TYPE_ACCELERATOR, &devices);
    return devices[0];
}

static const int DATA_SIZE = 1024;

int main(int argc, char *argv[])
{
    if(argc != 2)
    {
        std::cout << "Usage: " << argv[0] << " <XCLBIN File>" << std::endl;
        return EXIT_FAILURE;
    }

    std::string binary_file = argv[1];

    std::vector<int> source_a(DATA_SIZE, 10);
    std::vector<int> source_b(DATA_SIZE, 32);

    auto program_binary = read_binary_file(binary_file);
    cl::Program::Binaries bins{{program_binary.data(), program_binary.size()}};

    cl::Device device = select_device();
    cl::Context context(device, nullptr, nullptr, nullptr);
    cl::CommandQueue q(context, device, CL_QUEUE_PROFILING_ENABLE);

    cl::Program program(context, {device}, bins, nullptr);

    cl::Kernel vadd_kernel = cl::Kernel(program, "vector_add");

    cl::Buffer buffer_a(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, source_a.size() * sizeof(int), source_a.data());
    cl::Buffer buffer_b(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, source_b.size() * sizeof(int), source_b.data());
    cl::Buffer buffer_res(context, CL_MEM_READ_WRITE, source_a.size() * sizeof(int));

    int narg = 0;
    vadd_kernel.setArg(narg++, buffer_res);
    vadd_kernel.setArg(narg++, buffer_a);
    vadd_kernel.setArg(narg++, buffer_b);
    vadd_kernel.setArg(narg++, DATA_SIZE);

    q.enqueueTask(vadd_kernel);

    std::vector<int> result(DATA_SIZE, 0);
    q.enqueueReadBuffer(buffer_res, CL_TRUE, 0, result.size() * sizeof(int), result.data());

    int mismatch_count = 0;
    for(size_t i = 0; i < DATA_SIZE; ++i)
    {
        int host_result = source_a[i] + source_b[i];
        if(result[i] != host_result)
        {
            mismatch_count++;
            std::cout << "ERROR: " << result[i] << " != " << host_result << std::endl;
            break;
        }
    }

    std::cout << "RESULT: " << (mismatch_count == 0 ? "PASSED" : "FAILED") << std::endl;

    return 0;
}
```

The host-side code can now be compiled using GCC toolchain as:

```console
$ g++ host.cpp -I$XILINX_XRT/include -I$XILINX_VIVADO/include -lOpenCL -o host
```

The accelerator side (simple vector-add kernel) should be saved as `vadd.cl`.

```c++
#define BUFFER_SIZE 256
#define DATA_SIZE 1024

// TRIPCOUNT indentifier
__constant uint c_len = DATA_SIZE / BUFFER_SIZE;
__constant uint c_size = BUFFER_SIZE;

__attribute__((reqd_work_group_size(1, 1, 1)))
__kernel void vector_add(__global int* c,
    __global const int* a,
    __global const int* b,
    const int n_elements)
{
    int arrayA[BUFFER_SIZE];
    int arrayB[BUFFER_SIZE];

    __attribute__((xcl_loop_tripcount(c_len, c_len)))
    for (int i = 0; i < n_elements; i += BUFFER_SIZE)
    {
        int size = BUFFER_SIZE;

        if(i + size > n_elements)
            size = n_elements - i;

        __attribute__((xcl_loop_tripcount(c_size, c_size)))
        __attribute__((xcl_pipeline_loop(1))) readA:
        for(int j = 0; j < size; j++)
            arrayA[j] = a[i + j];

        __attribute__((xcl_loop_tripcount(c_size, c_size)))
        __attribute__((xcl_pipeline_loop(1))) readB:
        for(int j = 0; j < size; j++)
            arrayB[j] = b[i + j];

        __attribute__((xcl_loop_tripcount(c_size, c_size)))
        __attribute__((xcl_pipeline_loop(1))) vadd_writeC:
        for(int j = 0; j < size; j++)
            c[i + j] = arrayA[j] + arrayB[j];
    }
}
```

The accelerator-side code is build using Vitis `v++`.
This is three-step process, which either builds emulation binary or performs full HLS (depending on the value of `-t` argument).
The platform (specific accelerator) has to be also specified at this step (both for emulation and full HLS).

```console
$ v++ -c -t $IT4I_BUILD_MODE --platform $IT4I_PLATFORM -k vector_add -o vadd.xo vadd.cl
$ v++ -l -t $IT4I_BUILD_MODE --platform $IT4I_PLATFORM -o vadd.link.xclbin vadd.xo
$ v++ -p vadd.link.xclbin -t $IT4I_BUILD_MODE --platform $IT4I_PLATFORM -o vadd.xclbin
```

This process should result in `vadd.xclbin`, which can be loaded by host-side application.

### Running the Application

With both host application and kernel binary at hand the application (in emulation mode) can be launched as

```console
$ XCL_EMULATION_MODE=sw_emu ./host vadd.xclbin
```

or with real hardware (having compiled kernels with `IT4I_BUILD_MODE=hw`)

```console
./host vadd.xclbin
```

## Hybrid GPU and FPGA Application (HIP+XRT)

This simple 8-bit quantized dot product (`R = sum(X[i]*Y[i])`) example illustrates basic approach to utilize both GPU and FPGA accelerators in a single application.
The application takes the simplest approach, where both synchronization and data transfers are handled explicitly by the host.
The HIP toolchain is used to compile the single source host/GPU code as usual, but it is also linked with XRT runtime, which allows host to control the FPGA accelerator.
The FPGA kernels are built separately as in previous examples.

The host/GPU HIP code should be saved as `main.hip`

```c++
#include <iostream>
#include <vector>

#include "xrt/xrt_bo.h"
#include "experimental/xrt_xclbin.h"
#include "xrt/xrt_device.h"
#include "xrt/xrt_kernel.h"
#include "hip/hip_runtime.h"

const size_t DATA_SIZE = 1024;

float compute_reference(const float *srcX, const float *srcY, size_t count);

__global__ void quantize(int8_t *out, const float *in, size_t count)
{
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;

    for(size_t i = idx; i < count; i += blockDim.x * gridDim.x)
        out[i] = int8_t(in[i] * 127);
}

__global__ void dequantize(float *out, const int16_t *in, size_t count)
{
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;

    for(size_t i = idx; i < count; i += blockDim.x * gridDim.x)
        out[i] = float(in[i] / float(127*127));
}

int main(int argc, char *argv[])
{
    if(argc != 2)
    {
        std::cout << "Usage: " << argv[0] << " <XCLBIN File>" << std::endl;
        return EXIT_FAILURE;
    }

    // Prepare experiment data
    std::vector<float> srcX(DATA_SIZE);
    std::vector<float> srcY(DATA_SIZE);
    float outR = 0.0f;

    for(size_t i = 0; i < DATA_SIZE; ++i)
    {
        srcX[i] = float(rand()) / float(RAND_MAX);
        srcY[i] = float(rand()) / float(RAND_MAX);
        outR += srcX[i] * srcY[i];
    }

    float outR_quant = compute_reference(srcX.data(), srcY.data(), DATA_SIZE);

    std::cout << "REFERENCE: " << outR_quant << " (" << outR << ")" << std::endl;

    // Initialize XRT (FPGA device), load kernels binary and create kernel object
    xrt::device device(0);
    std::cout << "Loading xclbin file " << argv[1] << std::endl;
    xrt::uuid xclbinId = device.load_xclbin(argv[1]);
    xrt::kernel mulKernel(device, xclbinId, "multiply", xrt::kernel::cu_access_mode::exclusive);

    // Allocate GPU buffers
    float   *srcX_gpu, *srcY_gpu, *res_gpu;
    int8_t  *srcX_gpu_quant, *srcY_gpu_quant;
    int16_t *res_gpu_quant;
    hipMalloc(&srcX_gpu, DATA_SIZE * sizeof(float));
    hipMalloc(&srcY_gpu, DATA_SIZE * sizeof(float));
    hipMalloc(&res_gpu,  DATA_SIZE * sizeof(float));
    hipMalloc(&srcX_gpu_quant, DATA_SIZE * sizeof(int8_t));
    hipMalloc(&srcY_gpu_quant, DATA_SIZE * sizeof(int8_t));
    hipMalloc(&res_gpu_quant,  DATA_SIZE * sizeof(int16_t));

    // Allocate FPGA buffers
    xrt::bo srcX_fpga_quant(device, DATA_SIZE * sizeof(int8_t), mulKernel.group_id(0));
    xrt::bo srcY_fpga_quant(device, DATA_SIZE * sizeof(int8_t), mulKernel.group_id(1));
    xrt::bo res_fpga_quant(device, DATA_SIZE * sizeof(int16_t), mulKernel.group_id(2));

    // Copy experiment data from HOST to GPU
    hipMemcpy(srcX_gpu, srcX.data(), DATA_SIZE * sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(srcY_gpu, srcY.data(), DATA_SIZE * sizeof(float), hipMemcpyHostToDevice);

    // Execute quantization kernels on both input vectors
    quantize<<<16, 256>>>(srcX_gpu_quant, srcX_gpu, DATA_SIZE);
    quantize<<<16, 256>>>(srcY_gpu_quant, srcY_gpu, DATA_SIZE);

    // Map FPGA buffers into HOST memory, copy data from GPU to these mapped buffers and synchronize them into FPGA memory
    hipMemcpy(srcX_fpga_quant.map<int8_t *>(), srcX_gpu_quant, DATA_SIZE * sizeof(int8_t), hipMemcpyDeviceToHost);
    srcX_fpga_quant.sync(XCL_BO_SYNC_BO_TO_DEVICE);
    hipMemcpy(srcY_fpga_quant.map<int8_t *>(), srcY_gpu_quant, DATA_SIZE * sizeof(int8_t), hipMemcpyDeviceToHost);
    srcY_fpga_quant.sync(XCL_BO_SYNC_BO_TO_DEVICE);

    // Execute FPGA kernel (8-bit integer multiplication)
    auto kernelRun = mulKernel(res_fpga_quant, srcX_fpga_quant, srcY_fpga_quant, DATA_SIZE);
    kernelRun.wait();

    // Synchronize output FPGA buffer back to HOST and copy its contents to GPU buffer for dequantization
    res_fpga_quant.sync(XCL_BO_SYNC_BO_FROM_DEVICE);
    hipMemcpy(res_gpu_quant, res_fpga_quant.map<int16_t *>(), DATA_SIZE * sizeof(int16_t), hipMemcpyDeviceToHost);

    // Dequantize multiplication result on GPU
    dequantize<<<16, 256>>>(res_gpu, res_gpu_quant, DATA_SIZE);

    // Copy dequantized results from GPU to HOST
    std::vector<float> res(DATA_SIZE);
    hipMemcpy(res.data(), res_gpu, DATA_SIZE * sizeof(float), hipMemcpyDeviceToHost);

    // Perform simple sum on CPU
    float out = 0.0;
    for(size_t i = 0; i < DATA_SIZE; ++i)
        out += res[i];

    std::cout << "RESULT: " << out << std::endl;

    hipFree(srcX_gpu);
    hipFree(srcY_gpu);
    hipFree(res_gpu);
    hipFree(srcX_gpu_quant);
    hipFree(srcY_gpu_quant);
    hipFree(res_gpu_quant);

    return 0;
}

float compute_reference(const float *srcX, const float *srcY, size_t count)
{
    float out = 0.0f;

    for(size_t i = 0; i < count; ++i)
    {
        int16_t quantX(srcX[i] * 127);
        int16_t quantY(srcY[i] * 127);

        out += float(int16_t(quantX * quantY) / float(127*127));
    }

    return out;
}
```

The host/GPU application can be built using HIPCC as:

```console
$ hipcc -I$XILINX_XRT/include -I$XILINX_VIVADO/include -L$XILINX_XRT/lib -lxrt_coreutil main.hip -o host
```

The accelerator side (simple vector-multiply kernel) should be saved as `kernels.cpp`.

```c++
extern "C" {
    void multiply(
        short *out,
        const char *inX,
        const char *inY,
        int size)
    {
#pragma HLS INTERFACE m_axi port=inX bundle=aximm1
#pragma HLS INTERFACE m_axi port=inY bundle=aximm2
#pragma HLS INTERFACE m_axi port=out bundle=aximm1
        for(int i = 0; i < size; ++i)
            out[i] = short(inX[i]) * short(inY[i]);
    }
}
```

Once again the HLS kernel is build using Vitis `v++` in two steps:

```console
v++ -c -t $IT4I_BUILD_MODE --platform $IT4I_PLATFORM -k multiply kernels.cpp -o kernels.xo
v++ -l -t $IT4I_BUILD_MODE --platform $IT4I_PLATFORM kernels.xo -o kernels.xclbin
```

### Running the Application

In emulation mode (FPGA emulation, GPU HW is required) the application can be launched as:

```console
$ XCL_EMULATION_MODE=sw_emu ./host kernels.xclbin
REFERENCE: 256.554 (260.714)
Loading xclbin file ./kernels.xclbin
RESULT: 256.554
```

or, having compiled kernels with `IT4I_BUILD_MODE=hw` set, using real hardware (both FPGA and GPU HW is required)

```console
$ ./host kernels.xclbin
REFERENCE: 256.554 (260.714)
Loading xclbin file ./kernels.xclbin
RESULT: 256.554
```

## Additional Resources

- [https://xilinx.github.io/Vitis-Tutorials/][1]
- [http://xilinx.github.io/Vitis_Accel_Examples/][2]

[1]: https://xilinx.github.io/Vitis-Tutorials/
[2]: http://xilinx.github.io/Vitis_Accel_Examples/