Skip to content

Support buffer location for usm allocations #5661

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

Closed
Closed
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 @@ -7,7 +7,7 @@ NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are tradema

NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons.

This document describes an extension that adds a compile-time-constant property to tell the SYCL device compiler and runtime in which memory a particular accessor should be allocated. This is useful on targets that possess more than one type of global memory.
This document describes an extension that adds a compile-time-constant property to tell the SYCL device compiler and runtime in which memory a particular accessor or usm allocation should be allocated. This is useful on targets that possess more than one type of global memory.

== Name Strings

Expand Down Expand Up @@ -45,9 +45,31 @@ The use of this extension requires a target that supports SPV_INTEL_fpga_buffer_
== Overview

On targets that provide more than one type of global memory, knowing that a particular pointer can only access one of those memory types at compile time can enable compiler optimizations.
This extension adds an accessor property to indicate to the runtime in which of these memories the buffer corresponding to this accessor should be allocated and to inform the compiler that all accesses made through that accessor can only ever interact with the given memory.
This extension adds an accessor property to indicate to the runtime in which of these memories the buffer corresponding to this accessor should be allocated and to inform the compiler that all accesses made through that accessor can only ever interact with the given memory or it can be used to inform the usm alocation (`malloc_device`, `malloc_shared`, `malloc_host`) to allocate on a given memory.
This information is not a hint; it is a functional requirement of the program that must be respected.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you also add a section definition the feature-test macro for this extension? You can use the wording in the specification template in #5663. The name of the macro should be SYCL_EXT_INTEL_BUFFER_LOCATION.

Since there was no feature-test macro before, the value will always be 1. If any further changes are made to this extension, we'll bump the feature-test macro at that time.

Of course, the implementation in #5634 should also add a definition for this macro.

== Specification

=== Feature test macro

This extension provides a feature-test macro as described in the core SYCL
specification. An implementation supporting this extension must predefine the
macro `SYCL_INTEL_BUFFER_LOCATION` to one of the values defined in the table
below. Applications can test for the existence of this macro to determine if
the implementation supports this feature, or applications can test the macro's
value to determine which of the extension's features the implementation
supports.

[%header,cols="1,5"]
|===
|Value
|Description

|1
|Initial version of this extension.
|===


== Modifications to the SYCL 2020 Pre-Provisional Specification
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The description of the property (below) is currently very specific to accessors:

The buffer_location property notifies the SYCL device compiler that the given accessor will only ever point to the memory identified by the int template parameter of its instance class. It also notifies the SYCL runtime to store the given accessor in that memory.

Can you update this so it makes sense also for USM allocations?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the suggestion! The specifics of how this is going to interact with the malloc api depend on the which workaround we choose.

Given that we are only choosing between 1 and 3, malloc api will be updated to align with #5656

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks! I have added some more examples, and description for passing this property into usm allocation APIs.


=== Section 4.7.6.9.2 Device buffer accessor properties
Expand Down Expand Up @@ -86,6 +108,37 @@ It also notifies the SYCL runtime to store the given accessor in that memory. |
|===
--


== Buffer location as properties of USM allocation APIs

The buffer location property can also be passed to USM allocation APIs. The following is a synopsis.

[source,c++]
----
// A property list containing a compile-time property and a runtime property
sycl::ext::oneapi::experimental::properties properties{sycl::ext::intel::property::buffer_location<1>, some_runtime_prop(1)};
// The compile time property is passed to template arguments of annotated_ptr
auto data = sycl::ext::oneapi::experimental::malloc_device<int>(N, q, properties);
// data is of type annotated_ptr<int*, property_list_t<sycl::ext::intel::property::buffer_location<1>>>
sycl::queue q;
q.parallel_for(range<1>(N), [=] (id<1> i){
data[i] *= 2;
}).wait();
----

The table below describes the effects of associating the property with each malloc APIs.

|===
|USM API|Description

|`malloc_device`
|The returned device pointer should point to the target memory location.
|`malloc_shared`
|The returned pointer should be implicitly migrated to the target memory location.

|===


== Issues

== Revision History
Expand All @@ -96,4 +149,5 @@ It also notifies the SYCL runtime to store the given accessor in that memory. |
|========================================
|Rev|Date|Author|Changes
|1|2020-09-08|Joe Garvey|*Initial public draft*
|2|2022-03-03|Sherry Yuan|*USM API acceptance of buffer location*
|========================================