Support buffer location for usm allocations #5661
Conversation
| @@ -45,7 +45,7 @@ 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. | |||
|
|
|||
| == Modifications to the SYCL 2020 Pre-Provisional Specification | |||
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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
There was a problem hiding this comment.
Thanks! I have added some more examples, and description for passing this property into usm allocation APIs.
| @@ -45,7 +45,7 @@ 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. | |||
|
|
|||
There was a problem hiding this comment.
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.
|
I don't understand this PR. The existing extension "sycl_ext_intel_buffer_location" is a runtime property that can be passed to the |
The existing sycl_ext_intel_buffer_location is a compile time property that is passed to clang to generate correct SPIRV (At least this is what Dimitry told me). This is meant to complement the malloc+annotated pointer spec, as commented by Joe here; #5656 (comment) |
The existing Tagging @GarveyJoe |
|
sycl::ext::intel::property::buffer_location does in fact define a compile-time-constant property, but it builds on the older SYCL_ONEAPI_accessor_properties which was the precursor to sycl_ext_oneapi_properties. There's a column in table 4.50 labelled "Compile-time Constant" for which the value of this property is true. It can be passed to accessor_property_list, not sycl::property_list, and is preserved at compile time in the type of the accessor. Sherry, if you want to use this extension in combination with your new USM malloc properties extension then you'll need to update it to follow the approach from sycl_ext_oneapi_properties. The changes will be so significant, and the extension will be so bifurcated for the two different use cases, that it might be easier to create a new extension with a new property with a slightly different name. |
Right that make sense, in this case it is probably better to write a separate spec defining the property (which could possibly be directly in #5656 ), I will close this PR for this reason. Feel free to reopen it again if it turns out to be needed. |
The solution for supporting runtime & compile time properties for usm allocation requires buffer_locaiton property to be passed into the malloc APIs.
CC: @gmlueck @steffenlarsen