A newer version of this document is available. Customers should click here to go to the newest version.
Visible to Intel only — GUID: GUID-9F96D878-8F9D-4886-925A-7C92A77B51DA
Visible to Intel only — GUID: GUID-9F96D878-8F9D-4886-925A-7C92A77B51DA
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 component.
To customize the interface, use a functor to specify the component and use one of the two compiler-defined macros. The macros are defined in the sycl/ext/intel/prototype/interfaces.hpp header file.
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. For these interfaces, the host program must allocate the memory using the malloc_shared function that specifies 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::experimental::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) {
#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 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;
}