Skip to content

[SYCL][FPGA][NFC] LSU builtin test and documentation update #3546

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Apr 30, 2021
Merged

[SYCL][FPGA][NFC] LSU builtin test and documentation update #3546

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
520540f
[FPGA][NFC][LSU] Updating extension test and docs for LSU change to m…
Apr 14, 2021
288e730
[FPGA][NFC][LSU] Running clang-format on fpga.cpp lsu test
Apr 14, 2021
7b207f3
Merge branch 'sycl' of https://github.com/intel/llvm into doc_update
Apr 22, 2021
9b94312
Merge branch 'sycl' of https://github.com/intel/llvm into doc_update
Apr 26, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 9 additions & 5 deletions sycl/doc/extensions/IntelFPGA/FPGALsu.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@ is included in `CL/sycl/INTEL/fpga_extensions.hpp`.
The class `cl::sycl::INTEL::lsu` allows users to explicitly request that the
implementation of a global memory access is configured in a certain way. The
class has two member functions, `load()` and `store()` which allow loading from
and storing to a `global_ptr`, respectively, and is templated on the following
and storing to a `multi_ptr`, respectively, and is templated on the following
4 optional paremeters:

1. **`cl::sycl::INTEL::burst_coalesce<B>`, where `B` is a boolean**: request,
Expand Down Expand Up @@ -47,10 +47,12 @@ template <class... mem_access_params> class lsu final {
public:
lsu() = delete;

template <typename T> static T load(sycl::global_ptr<T> Ptr) {
template <typename _T, access::address_space _space>
static _T load(sycl::multi_ptr<_T, _space> Ptr) {
check_space<_space>();
check_load();
#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
return *__builtin_intel_fpga_mem((T *)Ptr,
return *__builtin_intel_fpga_mem((_T *)Ptr,
_burst_coalesce | _cache |
_dont_statically_coalesce | _prefetch,
_cache_val);
Expand All @@ -59,10 +61,12 @@ public:
#endif
}

template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
template <typename _T, access::address_space _space>
static void store(sycl::multi_ptr<_T, _space> Ptr, _T Val) {
check_space<_space>();
check_store();
#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
*__builtin_intel_fpga_mem((T *)Ptr,
*__builtin_intel_fpga_mem((_T *)Ptr,
_burst_coalesce | _cache |
_dont_statically_coalesce | _prefetch,
_cache_val) = Val;
Expand Down
95 changes: 64 additions & 31 deletions sycl/test/extensions/fpga.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,30 @@ template <unsigned ID> struct ethernet_pipe_id {
static constexpr unsigned id = ID;
};

template <typename T, cl::sycl::access::address_space space>
void lsu_body(cl::sycl::multi_ptr<T,space> input_ptr,cl::sycl::multi_ptr<T,space> output_ptr) {
using PrefetchingLSU =
cl::sycl::INTEL::lsu<cl::sycl::INTEL::prefetch<true>,
cl::sycl::INTEL::statically_coalesce<false>>;

using BurstCoalescedLSU =
cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>,
cl::sycl::INTEL::statically_coalesce<false>>;

using CachingLSU =
cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>,
cl::sycl::INTEL::cache<1024>,
cl::sycl::INTEL::statically_coalesce<false>>;

using PipelinedLSU = cl::sycl::INTEL::lsu<>;

int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]

BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
PipelinedLSU::store(output_ptr + 1, Y); // output_ptr[1] = Y
}

using ethernet_read_pipe =
sycl::INTEL::kernel_readable_io_pipe<ethernet_pipe_id<0>, int, 0>;
using ethernet_write_pipe =
Expand Down Expand Up @@ -58,39 +82,48 @@ int main() {

/*Check LSU interface*/
{
cl::sycl::buffer<int, 1> output_buffer(1);
auto *in_ptr = cl::sycl::malloc_host<int>(1, Queue.get_context());

Queue.submit([&](cl::sycl::handler &cgh) {
auto output_accessor =
output_buffer.get_access<cl::sycl::access::mode::write>(cgh);

cgh.single_task<class kernel>([=] {
cl::sycl::host_ptr<int> input_ptr(in_ptr);
auto output_ptr = output_accessor.get_pointer();

using PrefetchingLSU =
cl::sycl::INTEL::lsu<cl::sycl::INTEL::prefetch<true>,
cl::sycl::INTEL::statically_coalesce<false>>;

using BurstCoalescedLSU =
cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>,
cl::sycl::INTEL::statically_coalesce<false>>;

using CachingLSU =
cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>,
cl::sycl::INTEL::cache<1024>,
cl::sycl::INTEL::statically_coalesce<false>>;

using PipelinedLSU = cl::sycl::INTEL::lsu<>;

int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]

BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
PipelinedLSU::store(output_ptr + 1, Y); // output_ptr[1] = Y
{
auto *out_ptr = cl::sycl::malloc_host<int>(1, Queue.get_context());
auto *in_ptr = cl::sycl::malloc_host<int>(1, Queue.get_context());
Queue.submit([&](sycl::handler &cgh) {
cgh.single_task<class HostAnnotation>([=]() {
cl::sycl::host_ptr<int> input_ptr(in_ptr);
cl::sycl::host_ptr<int> output_ptr(out_ptr);
intelfpga::lsu_body<
int, cl::sycl::access::address_space::global_host_space>(
input_ptr, output_ptr);
});
});
});
}
{
auto *out_ptr = cl::sycl::malloc_device<int>(1, Queue);
auto *in_ptr = cl::sycl::malloc_device<int>(1, Queue);
Queue.submit([&](sycl::handler &cgh) {
cgh.single_task<class DeviceAnnotation>([=]() {
cl::sycl::device_ptr<int> input_ptr(in_ptr);
cl::sycl::device_ptr<int> output_ptr(out_ptr);
intelfpga::lsu_body<
int, cl::sycl::access::address_space::global_device_space>(
input_ptr, output_ptr);
});
});
}
{
cl::sycl::buffer<int, 1> output_buffer(1);
cl::sycl::buffer<int, 1> input_buffer(1);
Queue.submit([&](sycl::handler &cgh) {
auto output_accessor =
output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
auto input_accessor =
input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
cgh.single_task<class AccessorAnnotation>([=]() {
auto input_ptr = input_accessor.get_pointer();
auto output_ptr = output_accessor.get_pointer();
intelfpga::lsu_body<>(input_ptr, output_ptr);
});
});
}
}

return 0;
Expand Down