[SYCL] Add new kernel-arg-runtime-aligned metadata.

[srividya-sundaram]

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 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”.

[zjin-lcf]

I have some question. Are your changes related to #5113 ?
Thanks for describing the issue, cause, and solution.

[srividya-sundaram]

Not sure why SYCL / Default Linux / Build SYCL toolchain (pull_request) is failing . It seems to be stuck after the install step :
https://github.com/intel/llvm/runs/4479609101?check_suite_focus=true

[srividya-sundaram]

Fyi, I ran into some issues with my dev machine that caused delay in addressing the reviews comments. I am currently working on them.

[srividya-sundaram]

Ping @premanandrao @smanna12 @elizabethandrews

[smanna12]

LGTM

[bader]

@srividya-sundaram, please, fix "OCL x64 Test Suite" failures.

[srividya-sundaram]

@bader @alexbatashev I see this same error for all the failing tests in Linux / OCL x64 Test Suite.
I tried re-running the jobs a few times but still see the same errors.

Unsupported SPIR-V module
SPIRV module requires unsupported capability 5939
Compilation failed

https://github.com/intel/llvm/runs/4793458932?check_suite_focus=true#step:6:34113

Could you please comment on that?

[mohammadfawaz]

@bader @alexbatashev I see this same error for all the failing tests in Linux / OCL x64 Test Suite. I tried re-running the jobs a few times but still see the same errors.

Unsupported SPIR-V module
SPIRV module requires unsupported capability 5939
Compilation failed

https://github.com/intel/llvm/runs/4793458932?check_suite_focus=true#step:6:34113

Could you please comment on that?

That is odd - Capability is 5939 is ICapRuntimeAlignedAttributeINTEL which was introduced here: KhronosGroup/SPIRV-LLVM-Translator@974d4db and is meant to work well with this change.
@MrSidims any ideas what is going on here?

[bader]

I guess OpenCL CPU runtime doesn't support this extensions yet.

[MrSidims]

That is odd - Capability is 5939 is ICapRuntimeAlignedAttributeINTEL which was introduced here: KhronosGroup/SPIRV-LLVM-Translator@974d4db and is meant to work well with this change. @MrSidims any ideas what is going on here?

It an error from SPIR-V consumer. CPU runtime and FPGA emulator added a support for SPV_INTEL_runtime_aligned long ago, whilst GPU runtime is not yet here. AFAIU GPU can have little benefits from this metadata in some corner cases, but I'll disable the extension for it in a sec, leaving only FPGA H/W to support this extension.

[bader]

CPU runtime and FPGA emulator added a support for SPV_INTEL_runtime_aligned long ago,

From the logs:

terminate called after throwing an instance of 'cl::sycl::compile_program_error'
what(): The program was built for 1 devices
Build program log for 'Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz':
Compilation started
Unsupported SPIR-V module
SPIRV module requires unsupported capability 5939
Compilation failed
-11 (CL_BUILD_PROGRAM_FAILURE)

[MrSidims]

CPU runtime and FPGA emulator added a support for SPV_INTEL_runtime_aligned long ago,

From the logs:

terminate called after throwing an instance of 'cl::sycl::compile_program_error' what(): The program was built for 1 devices Build program log for 'Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz': Compilation started Unsupported SPIR-V module SPIRV module requires unsupported capability 5939 Compilation failed -11 (CL_BUILD_PROGRAM_FAILURE)

I see, then it just about CPU runtime version

Either way it should be save to just leave the extension be supported on FPGA H/W only for now.

[MrSidims]

I've created #5299 with this PR checkouted. Lets see if the pre-commit passes there.

[srividya-sundaram]

@MrSidims Hi, I now see the following error in all failing tests for SYCL / Default Linux / OCL x64 Test Suite

Build program log for 'AMD Ryzen 7 5800X 8-Core Processor             ':
Compilation started
Unsupported SPIR-V module
SPIRV module requires unsupported capability 5939
Compilation failed

And the following for SYCL / Default Linux / HIP AMD GPU Test Suite :

lld: error: undefined hidden symbol: __spirv_Group*(unsigned int, unsigned int, float)

[MrSidims]

> And the following for SYCL / Default Linux / HIP AMD GPU Test Suite :

lld: error: undefined hidden symbol: __spirv_Group*(unsigned int, unsigned int, float)

I believe it can not be related with neither your patch or SPIR-V extension.

Hi, I now see the following error in all failing tests for SYCL / Default Linux / OCL x64 Test Suite
Build program log for 'AMD Ryzen 7 5800X 8-Core Processor

Wow, didn't know, that we have testing on AMD for a regular basis.
The test failure you observe is unexpected. Code-wise rationale: the extension is now enabled only for FPGA H/W and hence the metadata this patch is adding will be ignored during translation to SPIR-V. CI rationale: you may see results of testing in #5299 (to where I initially checked out this PR), and there pre-commit is passing even on summary:run.
During some work in llvm-test-suite I observed, that restarting testing doesn't make CI to rebase PR (this issue was reported internally, unsure if it's fixed). May be you are facing something similar. If so, I'd suggest to retrigger CI jobs not by pressing 'restart' button, but by updating the PR itself.

[bader]

Tagging @zahiraam for awareness.