Load-Store Unit Controls
The
Intel® oneAPI
allows you to control the LSU that is generated for
global memory accesses via a set of templated options in the
DPC++/C++
CompilerINTEL::lsu
class. The
INTEL::lsu
class has two member functions,
load()
and
store()
, which allow loading from and storing to a global pointer, respectively.
Not all LSU control options listed below work with both
load()
and
store()
members function. Not all styles and types are supported on every target device. Thus, the use of LSU control options results in a specific type of LSU.
Control
| Syntax
| Value
| Default Value
| Supported Functionality
|
---|---|---|---|---|
Currently, not every combination of LSU controls is supported by the compiler. The following rules apply:
- For store, theINTEL::cachecontrol must be 0 and theINTEL::prefetchcontrol must be false.
- For load, if theINTEL::cachecontrol is set to a value greater than 0, then theINTEL::burst_coalescecontrol must be set to true.
- For load, exactly one ofINTEL::prefetchandINTEL::burst_coalescecontrol options is allowed to be true.
- For load, exactly one ofINTEL::prefetchandINTEL::cachecontrol options are allowed to be true.
Burst-Coalesce
The burst-coalesce control helps the
Intel® oneAPI
infer a burst-coalesced LSU. It can be combined with other attributes such as
DPC++/C++
Compilercache
. For more details about this LSU type, refer to
Burst-Coalesced Load-Store Units.
Consider the following example of burst-coalesce LSU control:
using BurstCoalescedLSU = cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>, cl::sycl::INTEL::statically_coalesce<false>>; BurstCoalescedLSU::store(output_ptr, X);
Static Coalescing
The static-coalescing control helps the
Intel® oneAPI
turn on or off static coalescing for global memory accesses. Static coalescing of memory accesses is the default behavior of the compiler, so this feature can be used to turn off static coalescing. It can be combined with other attributes such as
DPC++/C++
Compilerburst_coalesce
. For more details about this memory optimization modifier, refer to
Static Memory Coalescing.
Consider the following example using the control to turn off static-coalescing LSU control:
using NonStaticCoalescedLSU = cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>, cl::sycl::INTEL::statically_coalesce<false>>; int X = NonStaticCoalescedLSU::load(input_ptr);
Prefetch
The prefetch control helps the
Intel® oneAPI
infer a prefetch style LSU. The compiler honor this control only if it is physically possible. If there is a risk that an LSU is not functionally correct, it is not inferred. This control also works only for load accesses. For more details about this LSU type, refer to
Prefetching Load-Store Units.
DPC++/C++
CompilerConsider the following example of prefetch LSU control:
using PrefetchingLSU = cl::sycl::INTEL::lsu<cl::sycl::INTEL::prefetch<true>, cl::sycl::INTEL::statically_coalesce<false>>; int X = PrefetchingLSU::load(input_ptr);
Cache
The cache control helps the
Intel® oneAPI
turn on or off caching behavior and set the size of the cache. This LSU modifier can be useful for
DPC++/C++
Compilerparallel_for
kernels. The cache control can be combined with other attributes such as
burst_coalesce
. For more details about this LSU modifier, refer to
Cached.
Consider the following example of cache LSU control:
using CachingLSU = cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>, cl::sycl::INTEL::cache<1024>, cl::sycl::INTEL::statically_coalesce<false>>; int X = CachingLSU::load(input_ptr);
Pipeline
The pipeline control helps the
Intel® oneAPI
infer the LSU style of pipeline. It can be combined with other attributes such as
DPC++/C++
Compilerstatic coalesce
. For more details about this LSU type, refer to
Pipelined Load-Store Units.
For example:
using PipelinedLSU = cl::sycl::INTEL::lsu<>; PipelinedLSU::store(output_ptr,X);