Intel® oneAPI DPC++ FPGA Optimization Guide

Developer Guide

Contents

Load-Store Unit Controls

The
Intel® oneAPI
DPC++/C++
 Compiler
 allows you to control the LSU that is generated for global memory accesses via a set of templated options in the
INTEL::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.
LSU Control Options Available in the
INTEL::lsu
 Class
Control
Syntax
Value
Default Value
Supported Functionality
INTEL::burst_coalesce<B>
B
 is a boolean
False
Load and store
INTEL::statically_coalesce<B>
B
 is a boolean
True
Load and store
INTEL::prefetch<B>
B
 is a boolean
false
Only loads
INTEL::cache<N>
N
 is an integer greater than or equal to 0.
0
Only loads
No option provided
-
Defaults above result in this.
Load and store
Currently, not every combination of LSU controls is supported by the compiler. The following rules apply:
  • For store, the
    INTEL::cache
     control must be 0 and the
    INTEL::prefetch
     control must be false.
  • For load, if the
    INTEL::cache
     control is set to a value greater than 0, then the
    INTEL::burst_coalesce
     control must be set to true.
  • For load, exactly one of
    INTEL::prefetch
     and
    INTEL::burst_coalesce
     control options are allowed to be true.
  • For load, exactly one of
    INTEL::prefetch
     and
    INTEL::cache
     control options are allowed to be true.

Burst-Coalesce

The burst-coalesce control helps the
Intel® oneAPI
DPC++/C++
 Compiler
 infer a burst-coalesced LSU. It can be combined with other attributes such as
cache
. 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
DPC++/C++
 Compiler
 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
burst_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
DPC++/C++
 Compiler
 infer a prefetch style LSU. The compiler honors 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.
Consider 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
DPC++/C++
 Compiler
 turn on or off caching behavior and set the size of the cache. This LSU modifier can be useful for
parallel_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
DPC++/C++
 Compiler
 infer the LSU style of pipeline. It can be combined with other attributes such as
static 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);

Product and Performance Information

1

Performance varies by use, configuration and other factors. Learn more at www.Intel.com/PerformanceIndex.