Memory attributes for device_global

[v-m]

I am migrating IP component from HLS to oneAPI. The component is basic exponential moving average filtering.
Implementation has state, which is saved between invocation. In HLS this can be achieved using static variable inside function. Based on my research in oneAPI the recommended way is to to use device_global memory (see FPGA Optimization Guide for Intel®
oneAPI Toolkits - device_global Extension).

using FPGAProperties = decltype(sycl::ext::oneapi::experimental::properties(
  sycl::ext::oneapi::experimental::device_image_scope,
  sycl::ext::oneapi::experimental::host_access_none));

// Not able to apply memory attributes here
// [[intel::fpga_register]]
sycl::ext::oneapi::experimental::device_global<int, FPGAProperties> last;

struct ExpKernel {
  float sample;

  void operator()() const {
    last = 0.5f * last + 0.5f * sample;
    OutputPipe::write(last);
  }
};

This works as expected and generates device scoped memory. But I am not able to apply any memory attributes to this memory, like forcing registers or specifying bankwidth.

So my questions are:

1. How to apply memory attributes to device_global?

2. Is it the only way to achieve saving state across multiple invocations (besides using global memory)

[BoonBengT_Intel]

Hi @v-m,

Thank you for posting in Intel community forum, hope all is well and apologies for the delayed in response.

Noted and migration and interest in oneAPI, allow me to check this internally and will get back to you as soon as possible.

Best Wishes

BB

[BoonBengT_Intel]

Hi @v-m,

Thank you for the patients and apologies for the delayed in response.

If understood the situation correctly, you would like to implement the variable in the onchip memory as per the memory attributes.

For that the device_global variable is implmented in the onchip memory already during the hardware implementation.

You may find the details explanation in the section hardware implementation in the link below:

- https://www.intel.com/content/www/us/en/docs/oneapi-fpga-add-on/optimization-guide/2023-2/device-global-ext.html

Please do let us know if situation are misunderstood.

Best Wishes

BB

[v-m]

Hello

Thank you for your response. Yes, I need to implement the variable in the onchip memory and it can be achieved with device_global. But I am not able to apply any memory attributes to this device_global variable, for example [[intel::fpga_register]], [[intel::singlepump]], etc.
Can you please clarify if applying such attributes is possible in this case?

[BoonBengT_Intel]

Hi @v-m,

Noted on the intention, we are checking to confirm internally on the implementation if we can/should be done this way.

Will get back to you as soon as we have an updates.

Just to understand further, as mention you are not able to apply the memory attributes on device_global.

Is there any error or complain from the compiler that you are seeing?

Regards

BB

[v-m]

Hello

I've tried several things. I assumed those attributes must be applied to device_global declaration directly. Here is some tests:

// error: 'private_copies' attribute only applies to local non-const variables and non-static data members
[[intel::private_copies(2)]]
sycl::ext::oneapi::experimental::device_global<int, FPGAProperties> value;

// error: 'fpga_register' attribute only applies to constant variables, local variables, static variables, and non-static data members
[[intel::fpga_register]]
sycl::ext::oneapi::experimental::device_global<int, FPGAProperties> value;

// no errors, but based on report - attribute wasn't applied
[[intel::fpga_register]]
static sycl::ext::oneapi::experimental::device_global<int, FPGAProperties> value;

[BoonBengT_Intel]

Hi @v-m,

Thanks for holding and noted on the several tried from your end.

After some internal alignment and investigation, the question for can/should memory attributes be applied to device_global variable are not currently ready for oneAPI release 2024.0. It will be planned for the next release please stay tune.

The current device_global is implemented for memory enable with host access (i.e., write/read/none), but it does not let user define/configure yet.

Hope that clarify.

Regards

BB

[BoonBengT_Intel]

Hi @v-m,

Greetings, just checking in to see if there is any further doubts in regards to this matter.

Hope your doubts have been clarified.

Best Wishes

BB

[BoonBengT_Intel]

Hi @v-m,

Greetings, as we do not receive any further clarification/updates on the matter, hence would assume challenge are overcome. Please login to ‘https://supporttickets.intel.com’, view details of the desire request, and post a feed/response within the next 15 days to allow me to continue to support you. After 15 days, this thread will be transitioned to community support. For new queries, please feel free to open a new thread and we will be right with you. Pleasure having you here.

Best Wishes

BB