Skip to content

Commit 319ec1d

Browse files
intel-jisheng1intel-jisheng1
committed
sync
1 parent b5bdf36 commit 319ec1d

File tree

10 files changed

+1176
-0
lines changed

10 files changed

+1176
-0
lines changed

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/CMakeLists.txt

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
1+
if(UNIX)
2+
# Direct CMake to use icpx rather than the default C++ compiler/linker
3+
set(CMAKE_CXX_COMPILER icpx)
4+
else() # Windows
5+
# Force CMake to use icpx rather than the default C++ compiler/linker
6+
# (needed on Windows only)
7+
include (CMakeForceCompiler)
8+
CMAKE_FORCE_CXX_COMPILER (icx-cl IntelDPCPP)
9+
include (Platform/Windows-Clang)
10+
endif()
11+
12+
cmake_minimum_required (VERSION 3.4)
13+
14+
project(register_map_and_streaming_interfaces CXX)
15+
16+
set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
17+
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
18+
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
19+
20+
add_subdirectory (src)

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/README.md

Lines changed: 393 additions & 0 deletions
Large diffs are not rendered by default.

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/assets/register_map_invocation_interface.png

50.6 KB
Loading

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/assets/streaming_invocation_interface.png

32.9 KB
Loading

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/sample.json

Lines changed: 76 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,76 @@
1+
{
2+
"guid": "69415BED-D452-449A-8F5A-DB01ACCE38DC",
3+
"name": "invocation interfaces",
4+
"categories": ["Toolkit/oneAPI Direct Programming/C++SYCL FPGA/Tutorials/Features/experimental"],
5+
"description": "An Intel® FPGA tutorial demonstrating the usage of register_map and streaming invocation interfaces",
6+
"toolchain": ["icpx"],
7+
"os": ["linux", "windows"],
8+
"targetDevice": ["FPGA"],
9+
"builder": ["ide", "cmake"],
10+
"languages": [{"cpp":{}}],
11+
"commonFolder": {
12+
"base": "../../../..",
13+
"include": [
14+
"README.md",
15+
"Tutorials/Features/experimental/invocation_interfaces",
16+
"include"
17+
],
18+
"exclude": []
19+
},
20+
"ciTests": {
21+
"linux": [
22+
{
23+
"id": "fpga_emu",
24+
"steps": [
25+
"icpx --version",
26+
"mkdir build",
27+
"cd build",
28+
"cmake ..",
29+
"make fpga_emu",
30+
"./register_map_functor_model.fpga_emu",
31+
"./streaming_functor_model.fpga_emu",
32+
"./register_map_lambda_model.fpga_emu",
33+
"./streaming_lambda_model.fpga_emu"
34+
]
35+
},
36+
{
37+
"id": "report",
38+
"steps": [
39+
"icpx --version",
40+
"mkdir build",
41+
"cd build",
42+
"cmake ..",
43+
"make report"
44+
]
45+
}
46+
],
47+
"windows": [
48+
{
49+
"id": "fpga_emu",
50+
"steps": [
51+
"icpx --version",
52+
"cd ../../../..",
53+
"mkdir build",
54+
"cd build",
55+
"cmake -G \"NMake Makefiles\" ../Tutorials/Features/experimental/invocation_interfaces",
56+
"nmake fpga_emu",
57+
"register_map_functor_model.fpga_emu.exe",
58+
"streaming_functor_model.fpga_emu.exe",
59+
"register_map_lambda_model.fpga_emu.exe",
60+
"streaming_lambda_model.fpga_emu.exe"
61+
]
62+
},
63+
{
64+
"id": "report",
65+
"steps": [
66+
"icpx --version",
67+
"cd ../../../..",
68+
"mkdir build",
69+
"cd build",
70+
"cmake -G \"NMake Makefiles\" ../Tutorials/Features/experimental/invocatoin_interfaces",
71+
"nmake report"
72+
]
73+
}
74+
]
75+
}
76+
}

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/src/CMakeLists.txt

Lines changed: 212 additions & 0 deletions
Large diffs are not rendered by default.

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/src/register_map_functor_model.cpp

Lines changed: 121 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,121 @@
1+
#include <sycl/ext/intel/fpga_extensions.hpp>
2+
#include <sycl/ext/intel/prototype/interfaces.hpp>
3+
#include <sycl/sycl.hpp>
4+
5+
#include "exception_handler.hpp"
6+
7+
using ValueT = int;
8+
9+
// offloaded computation
10+
ValueT SomethingComplicated(ValueT val) { return (ValueT)(val * (val + 1)); }
11+
12+
/////////////////////////////////////////
13+
14+
struct FunctorRegisterMapIP {
15+
// Use the 'register_map' annotation on a kernel argument to specify it to be
16+
// a register map kernel argument.
17+
register_map ValueT *input;
18+
// Without the annotations, kernel arguments will be inferred to be register
19+
// map kernel arguments if the kernel invocation interface is register mapped,
20+
// and vise-versa.
21+
ValueT *output;
22+
// A kernel with a register map invocation interface can also independently
23+
// have streaming kernel arguments, when annotated by 'conduit'.
24+
conduit size_t n;
25+
register_map_interface void operator()() const {
26+
for (int i = 0; i < n; i++) {
27+
output[i] = SomethingComplicated(input[i]);
28+
}
29+
}
30+
};
31+
32+
int main(int argc, char *argv[]) {
33+
#if FPGA_SIMULATOR
34+
auto selector = sycl::ext::intel::fpga_simulator_selector_v;
35+
#elif FPGA_HARDWARE
36+
auto selector = sycl::ext::intel::fpga_selector_v;
37+
#else // #if FPGA_EMULATOR
38+
auto selector = sycl::ext::intel::fpga_emulator_selector_v;
39+
#endif
40+
41+
bool passed = true;
42+
43+
size_t count = 16;
44+
if (argc > 1) count = atoi(argv[1]);
45+
46+
if (count <= 0) {
47+
std::cerr << "ERROR: 'count' must be positive" << std::endl;
48+
return 1;
49+
}
50+
51+
try {
52+
// create the device queue
53+
sycl::queue q(selector, fpga_tools::exception_handler);
54+
55+
// make sure the device supports USM host allocations
56+
sycl::device d = q.get_device();
57+
58+
// Print out the device information.
59+
std::cout << "Running on device: "
60+
<< q.get_device().get_info<sycl::info::device::name>().c_str()
61+
<< std::endl;
62+
63+
if (!d.has(sycl::aspect::usm_host_allocations)) {
64+
std::cerr << "ERROR: The selected device does not support USM host"
65+
<< " allocations" << std::endl;
66+
return 1;
67+
}
68+
69+
ValueT *in = sycl::malloc_host<ValueT>(count, q);
70+
ValueT *functor_register_map_out = sycl::malloc_host<ValueT>(count, q);
71+
ValueT *golden = sycl::malloc_host<ValueT>(count, q);
72+
73+
// create input and golden output data
74+
for (int i = 0; i < count; i++) {
75+
in[i] = rand() % 77;
76+
golden[i] = SomethingComplicated(in[i]);
77+
functor_register_map_out[i] = 0;
78+
}
79+
80+
// validation lambda
81+
auto validate = [](auto &in, auto &out, size_t size) {
82+
for (int i = 0; i < size; i++) {
83+
if (out[i] != in[i]) {
84+
std::cout << "out[" << i << "] != in[" << i << "]"
85+
<< " (" << out[i] << " != " << in[i] << ")" << std::endl;
86+
return false;
87+
}
88+
}
89+
return true;
90+
};
91+
92+
// Launch the kernel with a register map invocation interface implemented in
93+
// the functor programming model
94+
std::cout << "Running the kernel with a register map invocation interface "
95+
"implemented in "
96+
"the functor programming model"
97+
<< std::endl;
98+
q.single_task(FunctorRegisterMapIP{in, functor_register_map_out, count})
99+
.wait();
100+
std::cout << "\t Done" << std::endl;
101+
102+
passed &= validate(golden, functor_register_map_out, count);
103+
std::cout << std::endl;
104+
105+
sycl::free(in, q);
106+
sycl::free(functor_register_map_out, q);
107+
sycl::free(golden, q);
108+
} catch (sycl::exception const &e) {
109+
// Catches exceptions in the host code
110+
std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n";
111+
std::terminate();
112+
}
113+
114+
if (passed) {
115+
std::cout << "PASSED\n";
116+
return 0;
117+
} else {
118+
std::cout << "FAILED\n";
119+
return 1;
120+
}
121+
}

‎DirectProgramming/C++SYCL_FPGA/Tutorials/Features/experimental/invocation_interfaces/src/register_map_lambda_model.cpp

Lines changed: 117 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,117 @@
1+
#include <sycl/ext/intel/fpga_extensions.hpp>
2+
#include <sycl/ext/intel/prototype/interfaces.hpp>
3+
#include <sycl/sycl.hpp>
4+
5+
#include "exception_handler.hpp"
6+
7+
using ValueT = int;
8+
// Forward declare the kernel names in the global scope.
9+
// This FPGA best practice reduces name mangling in the optimization reports.
10+
class LambdaRegisterMapIP;
11+
12+
// offloaded computation
13+
ValueT SomethingComplicated(ValueT val) { return (ValueT)(val * (val + 1)); }
14+
15+
/////////////////////////////////////////
16+
17+
void TestLambdaRegisterMapKernel(sycl::queue &q, ValueT *in, ValueT *out,
18+
size_t count) {
19+
// In the Lambda programming model, all kernel arguments will have the same
20+
// interface as the kernel invocation interface.
21+
q.single_task<LambdaRegisterMapIP>([=] register_map_interface {
22+
for (int i = 0; i < count; i++) {
23+
out[i] = SomethingComplicated(in[i]);
24+
}
25+
})
26+
.wait();
27+
28+
std::cout << "\t Done" << std::endl;
29+
}
30+
31+
int main(int argc, char *argv[]) {
32+
#if FPGA_SIMULATOR
33+
auto selector = sycl::ext::intel::fpga_simulator_selector_v;
34+
#elif FPGA_HARDWARE
35+
auto selector = sycl::ext::intel::fpga_selector_v;
36+
#else // #if FPGA_EMULATOR
37+
auto selector = sycl::ext::intel::fpga_emulator_selector_v;
38+
#endif
39+
40+
bool passed = true;
41+
42+
size_t count = 16;
43+
if (argc > 1) count = atoi(argv[1]);
44+
45+
if (count <= 0) {
46+
std::cerr << "ERROR: 'count' must be positive" << std::endl;
47+
return 1;
48+
}
49+
50+
try {
51+
// create the device queue
52+
sycl::queue q(selector, fpga_tools::exception_handler);
53+
54+
// make sure the device supports USM host allocations
55+
sycl::device d = q.get_device();
56+
57+
// Print out the device information.
58+
std::cout << "Running on device: "
59+
<< q.get_device().get_info<sycl::info::device::name>().c_str()
60+
<< std::endl;
61+
62+
if (!d.has(sycl::aspect::usm_host_allocations)) {
63+
std::cerr << "ERROR: The selected device does not support USM host"
64+
<< " allocations" << std::endl;
65+
return 1;
66+
}
67+
68+
ValueT *in = sycl::malloc_host<ValueT>(count, q);
69+
ValueT *lambda_register_map_out = sycl::malloc_host<ValueT>(count, q);
70+
ValueT *golden = sycl::malloc_host<ValueT>(count, q);
71+
72+
// create input and golden output data
73+
for (int i = 0; i < count; i++) {
74+
in[i] = rand() % 77;
75+
golden[i] = SomethingComplicated(in[i]);
76+
lambda_register_map_out[i] = 0;
77+
}
78+
79+
// validation lambda
80+
auto validate = [](auto &in, auto &out, size_t size) {
81+
for (int i = 0; i < size; i++) {
82+
if (out[i] != in[i]) {
83+
std::cout << "out[" << i << "] != in[" << i << "]"
84+
<< " (" << out[i] << " != " << in[i] << ")" << std::endl;
85+
return false;
86+
}
87+
}
88+
return true;
89+
};
90+
91+
// Launch the kernel with a register map invocation interface implemented in
92+
// the lambda programming model
93+
std::cout << "Running kernel with a register map invocation interface "
94+
"implemented in the "
95+
"lambda programming model"
96+
<< std::endl;
97+
TestLambdaRegisterMapKernel(q, in, lambda_register_map_out, count);
98+
passed &= validate(golden, lambda_register_map_out, count);
99+
std::cout << std::endl;
100+
101+
sycl::free(in, q);
102+
sycl::free(lambda_register_map_out, q);
103+
sycl::free(golden, q);
104+
} catch (sycl::exception const &e) {
105+
// Catches exceptions in the host code
106+
std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n";
107+
std::terminate();
108+
}
109+
110+
if (passed) {
111+
std::cout << "PASSED\n";
112+
return 0;
113+
} else {
114+
std::cout << "FAILED\n";
115+
return 1;
116+
}
117+
}

0 commit comments

Comments
 (0)