[SYCL] implement no_offset property for accessor_property_list class #4920
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@intel/llvm-reviewers-runtime friendly ping |
vladimirlaz
reviewed
Nov 11, 2021
Co-authored-by: vladimirlaz <[email protected]>
vladimirlaz
previously approved these changes
Nov 11, 2021
Co-authored-by: Steffen Larsen <[email protected]>
HabKaffee
dismissed stale reviews from steffenlarsen and vladimirlaz
via
a1cef53
vladimirlaz
approved these changes
Nov 12, 2021
|
@HabKaffee, it looks like post-commit build has failed? Could you please prepare fix ASAP? |
|
@vladimirlaz, yes, sure, I'm working on it |
MrSidims
reviewed
Nov 12, 2021
| // RUN: %clangxx -fsycl-device-only -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -S -emit-llvm -o - %s | FileCheck %s | ||
|
|
||
| #include <CL/sycl.hpp> | ||
|
|
There was a problem hiding this comment.
It would be nice to have E2E test in https://github.com/intel/llvm-test-suite as well. Are we planning to add it there? Basic requirement for the test to check:
- It doesn't crash during AOT and JIT compilation and execution on different devices (CPU/GPU/FPGA) (I assume, later the IR will be translated to SPIR-V and we need to check that this SPIR-V is acceptable for all of the mentioned devices (I'm sure, that it will be accepted, but still));
- Basic math operations with 'no-offsetted' accessors are having the correct result.
bader
pushed a commit
that referenced
this pull request
Jan 18, 2022
In this PR, if a kernel pointer argument comes from a global accessor, we generate a new metadata(kernel_arg_runtime_aligned) to the kernel to indicate that this pointer has runtime allocated alignment. If this information is available to the FPGA backend and if the accessor has no offset (e.g. through the user applying the [no_offset](#4920) property to their accessor), improvements to area of loads and stores can be made by using aligned LSUs. Without this enhancement we will continue to experience excess area. The SYCL spec already guarantees that accessors are aligned to some runtime-specific alignment. So the user's source doesn't need to change to provide the backend with this guarantee, we simply need to allow this information to propagate to the backend. Current IR implementation for kernel pointer argument from accessor looks like: `define dso_local spir_kernel void @_ZTSZ4mainE8kernel_A(i32 addrspace(1)* %_arg_, %"struct.cl::sycl::range"* byval(%"struct.cl::sycl::range") align 4 %_arg_1, %"struct.cl::sycl::range"* byval(%"struct.cl::sycl::range") align 4 %_arg_2, %"struct.cl::sycl::id"* byval(%"struct.cl::sycl::id") align 4 %_arg_3) #0 !kernel_arg_buffer_location !4 {` The new implementation will look like: `define dso_local spir_kernel void @_ZTSZ4mainE8kernel_A(i32 addrspace(1)* %_arg_, %"struct.cl::sycl::range"* byval(%"struct.cl::sycl::range") align 4 %_arg_1, %"struct.cl::sycl::range"* byval(%"struct.cl::sycl::range") align 4 %_arg_2, %"struct.cl::sycl::id"* byval(%"struct.cl::sycl::id") align 4 %_arg_3) #0 !kernel_arg_buffer_location !4 !kernel_arg_runtime_aligned !5 {` `!5 = !{i1 true, i1 false, i1 false, i1 false}` The metadata is applied to the kernel but really carries data about the kernel’s arguments. The first element of the metadata maps to the first kernel argument, the second to the second and so on. For this particular metadata the request is that the value of any metadata element is 'true' for any kernel arguments that corresponds to the base pointer of an accessor and 'false' otherwise. Accessors are handled specially by the frontend (because they are marked with sycl_special_class) and when a user captures an accessor in their SYCL kernel the FE splits up the single accessor into 4 separate kernel arguments. The first of those 4 arguments is a pointer and is the base pointer of the accessor. That pointer is known to have runtime-specific alignment and thus the element of the kernel-arg-runtime metadata that corresponds to that argument will have a value of “true”.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
no_offset property allows to reduce the amount of kernel arguments
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/accessor_properties/SYCL_ONEAPI_accessor_properties.asciidoc