A newer version of this document is available. Customers should click here to go to the newest version.
Visible to Intel only — GUID: GUID-EDDB340C-3B1D-4DE9-8DFE-E2E183F754D1
Visible to Intel only — GUID: GUID-EDDB340C-3B1D-4DE9-8DFE-E2E183F754D1
Memory-Mapped Interface Using Unified Shared Memory
You can use unified shared memory (USM) to fully customize a memory-mapped interface when compiling for an IP-only flow.
To customize the interface, use a functor to specify the component and use one of the two compiler-defined macros.
The following macro creates a memory-mapped host interface with the specified parameters. The base pointer is passed in through the register map.
register_map_mmhost( BL1, // buffer_location or aspace 28, // address width 64, // data width 16, // ! latency, must be atleast 16 0, // read_write_mode, 0: ReadWrite, 1: Read, 2: Write 1, // maxburst 0, // align, 0 defaults to alignment of the type 1 // waitrequest, 0: false, 1: true ) int *x;
You can also use the following macro instead to have the base pointer passed in through a conduit interface.
conduit_mmhost( BL1, // buffer_location or aspace 28, // address width 64, // data width 16, // ! latency, must be atleast 16 0, // read_write_mode, 0: ReadWrite, 1: Read, 2: Write 1, // maxburst 0, // align, 0 defaults to alignment of the type 1 // waitrequest, 0: false, 1: true ) int *x;
When you specify the macro properties, the order of the properties must be preserved. The compiler exits with an error out when you provide an unsupported combination. You can customize the following properties:
Property |
Description |
|---|---|
Buffer Location |
Specify the interface ID, which allows you to create multiple different interfaces. Buffer locations must be sequential integers, starting with 0. |
Address Width |
Width of the address bus |
Data Width |
Width of the data bus |
Latency |
Latency of the memory. Use 0 to specify a variable latency memory. |
Readwrite mode |
0: ReadWrite 1: Read 2: Write |
Max burst |
Set the maximum burst size. |
Align |
Memory alignment. 0 defaults to the type alignment. |
Wait Request |
Enable wait request: 0: false 1: true |
To include this macro in your program, create a kernel as a functor, allocate the memory on the host, and copy the data from the host to the kernel and back.
The following example creates two memory-mapped interfaces. The host program must allocate the memory using malloc_shared. Also, this allocation requires the buffer location as a property.
It initializes two values to 0. The kernel code then sets them to 5 and 6, respectively. It copies the desired memory locations to the host program, frees the allocated memory, and then verifies the output is as expected.
#include <sycl/sycl.hpp>
#include <sycl/ext/intel/fpga_extensions.hpp>
#include <sycl/ext/intel/prototype/interfaces.hpp>
using namespace sycl;
using ext::intel::prototype::property::usm::buffer_location;
constexpr int BL1 = 0;
constexpr int BL2 = 1;
struct MyIP {
register_map_mmhost(
BL1, // buffer_location or aspace
28, // address width
64, // data width
16, // ! latency, must be at least 16
0, // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
1, // maxburst
0, // align, 0 defaults to alignment of the type
1 // waitrequest, 0: false, 1: true
) int *x;
register_map_mmhost(
BL2, // buffer_location or aspace
28, // address width
64, // data width
16, // ! latency, must be at least 16
0, // read_write_mode, 0: ReadWrite, 1: Read, 2: Write
1, // maxburst
0, // align, 0 defaults to alignment of the type
1 // waitrequest, 0: false, 1: true
) int *y;
MyIP(int *x_, int *y_)
: x(x_), y(y_) {}
register_map_interface
void operator()() const {
*x = 5;
*y = 6;
}
};
void Test(int *first, int *second) {
#ifdef FPGA_EMULATOR
sycl::ext::intel::fpga_emulator_selector my_selector;
#elif FPGA_SIMULATOR
sycl::ext::intel::fpga_simulator_selector my_selector;
#else
sycl::ext::intel::fpga_selector my_selector;
#endif
queue q(my_selector);
int *HostA = malloc_shared<int>(sizeof(int), q, property_list{buffer_location(BL1)});
*HostA = 0;
int *HostB = malloc_shared<int>(sizeof(int), q, property_list{buffer_location(BL2)});
*HostB = 0;
q.single_task(MyIP{HostA, HostB}).wait();
*first = *HostB;
*second = *HostA;
sycl::free(HostA, q);
sycl::free(HostB, q);
}
int main() {
int first = 0;
int second = 0;
Test(&first, &second);
if (first == 6 && second == 5) std::cout << "PASSED\n";
else std::cout << "FAILED\n";
return 0;
}