diff --git a/sycl/doc/extensions/IntelFPGA/FPGALsu.md b/sycl/doc/extensions/IntelFPGA/FPGALsu.md index 7a412128d9afc..6bb10604a634d 100644 --- a/sycl/doc/extensions/IntelFPGA/FPGALsu.md +++ b/sycl/doc/extensions/IntelFPGA/FPGALsu.md @@ -1,41 +1,41 @@ # FPGA lsu -The Intel FPGA `lsu` class is implemented in `CL/sycl/intel/fpga_lsu.hpp` which -is included in `CL/sycl/intel/fpga_extensions.hpp`. +The Intel FPGA `lsu` class is implemented in `CL/sycl/INTEL/fpga_lsu.hpp` which +is included in `CL/sycl/INTEL/fpga_extensions.hpp`. -The class `cl::sycl::intel::lsu` allows users to explicitly request that the +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 4 optional paremeters: -1. **`cl::sycl::intel::burst_coalesce`, where `B` is a boolean**: request, +1. **`cl::sycl::INTEL::burst_coalesce`, where `B` is a boolean**: request, to the extent possible, that a dynamic burst coalescer be implemented when `load` or `store` are called. The default value of this parameter is `false`. -2. **`cl::sycl::intel::cache`, where `N` is an integer greater or equal to +2. **`cl::sycl::INTEL::cache`, where `N` is an integer greater or equal to 0**: request, to the extent possible, that a read-only cache of the specified size in bytes be implemented when when `load` is called. It is not allowed to use that parameter for `store`. The default value of this parameter is `0`. -3. **`cl::sycl::intel::statically_coalesce`, where `B` is a boolean**: +3. **`cl::sycl::INTEL::statically_coalesce`, where `B` is a boolean**: request, to the extent possible, that `load` or `store` accesses, is allowed to be statically coalesced with other memory accesses at compile time. The default value of this parameter is `true`. -4. **`cl::sycl::intel::prefetch`, where `N` is a boolean**: request, to the +4. **`cl::sycl::INTEL::prefetch`, where `B` is a boolean**: request, to the extent possible, that a prefetcher be implemented when `load` is called. It is not allowed to use that parameter for `store`. The default value of this parameter is `false`. Currently, not every combination of parameters is allowed due to limitations in the backend. The following rules apply: -1. For `store`, `cl::sycl::intel::cache` must be `0` and -`cl::sycl::intel::prefetch` must be `false`. -2. For `load`, if `cl::sycl::intel::cache` is set to a value greater than `0`, -then `cl::sycl::intel::burst_coalesce` must be set to `true`. -3. For `load`, exactly one of `cl::sycl::intel::prefetch` and -`cl::sycl::intel::burst_coalesce` is allowed to be `true`. -4. For `load`, exactly one of `cl::sycl::intel::prefetch` and -`cl::sycl::intel::cache` is allowed to be `true`. +1. For `store`, `cl::sycl::INTEL::cache` must be `0` and +`cl::sycl::INTEL::prefetch` must be `false`. +2. For `load`, if `cl::sycl::INTEL::cache` is set to a value greater than `0`, +then `cl::sycl::INTEL::burst_coalesce` must be set to `true`. +3. For `load`, exactly one of `cl::sycl::INTEL::prefetch` and +`cl::sycl::INTEL::burst_coalesce` is allowed to be `true`. +4. For `load`, exactly one of `cl::sycl::INTEL::prefetch` and +`cl::sycl::INTEL::cache` is allowed to be `true`. ## Implementation @@ -47,7 +47,7 @@ template class lsu final { public: lsu() = delete; - template static T &load(sycl::global_ptr Ptr) { + template static T load(sycl::global_ptr Ptr) { check_load(); #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem) return *__builtin_intel_fpga_mem((T *)Ptr, @@ -77,7 +77,7 @@ public: ## Usage ```c++ -#include +#include ... cl::sycl::buffer output_buffer(output_data, 1); cl::sycl::buffer input_buffer(input_data, 1); @@ -91,19 +91,19 @@ Queue.submit([&](cl::sycl::handler &cgh) { auto output_ptr = output_accessor.get_pointer(); using PrefetchingLSU = - cl::sycl::intel::lsu, - cl::sycl::intel::statically_coalesce>; + cl::sycl::INTEL::lsu, + cl::sycl::INTEL::statically_coalesce>; using BurstCoalescedLSU = - cl::sycl::intel::lsu, - cl::sycl::intel::statically_coalesce>; + cl::sycl::INTEL::lsu, + cl::sycl::INTEL::statically_coalesce>; using CachingLSU = - cl::sycl::intel::lsu, - cl::sycl::intel::cache<1024>, - cl::sycl::intel::statically_coalesce>; + cl::sycl::INTEL::lsu, + cl::sycl::INTEL::cache<1024>, + cl::sycl::INTEL::statically_coalesce>; - using PipelinedLSU = cl::sycl::intel::lsu<>; + 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] diff --git a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp index a52723c0c4a2e..4377ceacffcda 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp @@ -1,4 +1,4 @@ -//==-------------- fpga_lsu.hpp --- SYCL FPGA Reg Extensions ---------------==// +//==-------------- fpga_lsu.hpp --- SYCL FPGA LSU Extensions ---------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -47,7 +47,7 @@ template class lsu final { public: lsu() = delete; - template static T &load(sycl::global_ptr Ptr) { + template static T load(sycl::global_ptr Ptr) { check_load(); #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem) return *__builtin_intel_fpga_mem((T *)Ptr,