Skip to content

Adapt to the get_pointer return type change from upstream intel/llvm #1812

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.

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 3 commits into from
Aug 8, 2023
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
Original file line number Diff line number Diff line change
Expand Up @@ -45,8 +45,10 @@ void KernelRun(const std::vector<int> &in_data, std::vector<int> &out_data,
sycl::no_init);

h.single_task<LatencyControl>([=]() [[intel::kernel_args_restrict]] {
auto in_ptr = in_accessor.get_pointer();
auto out_ptr = out_accessor.get_pointer();
auto in_ptr =
in_accessor.template get_multi_ptr<sycl::access::decorated::no>();
auto out_ptr =
out_accessor.template get_multi_ptr<sycl::access::decorated::no>();

for (size_t i = 0; i < size; i++) {
// The following load has a label 0.
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@ This sample is an FPGA tutorial that demonstrates how to configure the load-stor

## Purpose

The compiler creates load-store units (LSU) to access off-chip data. The compiler has many options to choose from when configuring each LSU. The SYCL*-compliant LSU controls extension allows you to override the compiler's internal heuristics and control the architecture of each LSU. An introduction to the extension in this tutorial will explain the available options, extension defaults, appropriate use cases, and area trade-offs.
The compiler creates load-store units (LSU) to access memories, both on-chip and off-chip. The compiler has many options to choose from when configuring each LSU. The SYCL*-compliant LSU controls extension allows you to override the compiler's internal heuristics and control the architecture of individual LSUs that are used to access variable-latency off-chip memory. An introduction to the extension in this tutorial will explain the available options, extension defaults, appropriate use cases, and area trade-offs.

## Prerequisites

Expand Down Expand Up @@ -63,7 +63,7 @@ The sample illustrates the following important concepts.

### LSUs and LSU Styles

An LSU is a block that handles loading and storing data to and from memory. Off-chip memory can have variable latency. To mitigate this, different LSU implementations, referred to as styles, are available.
An LSU is a block that handles loading and storing data to and from memory. Off-chip memory can have variable latency. To mitigate this, different LSU styles are available.

The two LSU styles used in this tutorial are listed below:

Expand All @@ -77,11 +77,15 @@ The best LSU style depends on the memory access pattern in your design. There ar

In addition to these two styles, there are also LSU modifiers. LSU modifiers are add-ons that can be combined with LSU styles, such as caching, which can be combined with the burst-coalesced LSU style.

For more details on LSU modifiers and LSU styles, refer to the Memory Accesses section in the [FPGA Optimization Guide for Intel® oneAPI Toolkits Developer Guide](https://software.intel.com/content/www/us/en/develop/documentation/oneapi-fpga-optimization-guide).
For more details on LSU modifiers and LSU styles, refer to the *Memory Accesses* section in the [FPGA Optimization Guide for Intel® oneAPI Toolkits Developer Guide](https://www.intel.com/content/www/us/en/docs/oneapi-fpga-add-on/optimization-guide/current/memory-accesses.html).

### Introduction to the LSU Control Extension

The class: ```ext::intel::lsu``` enables you to control the architecture of the LSU. The class has two member functions, `load()` and `store()`, which allow loading from and storing to a global pointer.
The class: ```ext::intel::lsu``` enables you to control the architecture of the LSU. The class has two member functions, `load()` and `store()`, which allow loading from and storing to a global pointer (via `sycl::multi_ptr` rather than raw pointer).

There are two steps to use the LSU control extension to optimize LSU behaviour:
1. Get a `sycl::multi_ptr` representation of the memory you wish to access using the `get_multi_ptr<>()` function.
2. Access this `sycl::multi_ptr` using one of the LSU control functions.

The table below summarizes the LSU control extension parameters. The parameters will be respected to the extent possible.

Expand All @@ -97,17 +101,17 @@ If the default options are used, a pipelined LSU is implemented.
#### Example: Controlling the `prefetch` and `statically_coalesce` Parameters

```c++
//Creating typedefs using the LSU controls class
//for each combination of LSU options desired.
// Creating typedefs using the LSU controls class
// for each combination of LSU options desired.
using PrefetchingLSU = ext::intel::lsu<ext::intel::prefetch<true>,
ext::intel::statically_coalesce<false>>;
ext::intel::statically_coalesce<false>>;
// ...
q.submit([&](handler &h) {
h.single_task<Kernel>([=] {
//Pointer to external memory
auto input_ptr = input_accessor.get_pointer();
// Pointer to external memory
auto input_ptr = input_accessor.template get_multi_ptr<access::decorated::no>();

//Compiler will use a Prefetch LSU for this load
// Compiler will use a Prefetch LSU for this load
int in_data = PrefetchingLSU::load(input_ptr);

//...
Expand All @@ -116,7 +120,7 @@ q.submit([&](handler &h) {
```

Currently, not every combination of parameters is valid in the compiler.
For more details on the descriptions of LSU controls, styles, and modifiers refer to the *FPGA LSU Controls* section in the [FPGA Optimization Guide for Intel® oneAPI Toolkits Developer Guide](https://software.intel.com/content/www/us/en/develop/documentation/oneapi-fpga-optimization-guide).
For more details on the descriptions of LSU controls, styles, and modifiers refer to the *Load-Store Unit Controls* section in the [FPGA Optimization Guide for Intel® oneAPI Toolkits Developer Guide](https://www.intel.com/content/www/us/en/docs/oneapi-fpga-add-on/optimization-guide/current/load-store-unit-controls.html).

### Tutorial Overview

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -79,7 +79,8 @@ void KernelRun(const std::vector<int> &input_data, const size_t &input_size,

// Kernel that uses the prefetch LSU
h.single_task<KernelPrefetch>([=]() [[intel::kernel_args_restrict]] {
auto input_ptr = input_a.get_pointer();
auto input_ptr =
input_a.template get_multi_ptr<access::decorated::no>();
auto output_ptr = output_a.get_pointer();

int total = 0;
Expand All @@ -96,7 +97,8 @@ void KernelRun(const std::vector<int> &input_data, const size_t &input_size,

// Kernel that uses the burst-coalesced LSU
h.single_task<KernelBurst>([=]() [[intel::kernel_args_restrict]] {
auto input_ptr = input_a.get_pointer();
auto input_ptr =
input_a.template get_multi_ptr<access::decorated::no>();
auto output_ptr = output_a.get_pointer();

int total = 0;
Expand Down