diff options
author | Alex Voicu <alexandru.voicu@amd.com> | 2025-01-07 12:01:31 +0200 |
---|---|---|
committer | GitHub <noreply@github.com> | 2025-01-07 12:01:31 +0200 |
commit | 66acb2694655321b37a1ee3ff19207a111756562 (patch) | |
tree | 10cea448f00ce15570fae5ef4ff11d58cfdf021c /clang/lib/Frontend/CompilerInvocation.cpp | |
parent | 7ce15f3ba76e1780935715cfdd9dd368ebd23163 (diff) | |
download | llvm-66acb2694655321b37a1ee3ff19207a111756562.zip llvm-66acb2694655321b37a1ee3ff19207a111756562.tar.gz llvm-66acb2694655321b37a1ee3ff19207a111756562.tar.bz2 |
[clang][CodeGen][SPIRV] Translate `amdgpu_flat_work_group_size` into `max_work_group_size`. (#116820)
HIPAMD relies on the `amdgpu_flat_work_group_size` attribute to
implement key functionality such as the `__launch_bounds__` `__global__`
function annotation. This attribute is not available / directly
translatable to SPIR-V, hence as it is AMDGCN flavoured SPIR-V suffers
from information loss.
This patch addresses that limitation by converting the unsupported
attribute into the `max_work_group_size` attribute which maps to
[`MaxWorkgroupSizeINTEL`](https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_kernel_attributes.asciidoc),
which is available in / handled by SPIR-V. When reverse translating from
SPIR-V to AMDGCN LLVMIR we invert the map and add the original AMDGPU
attribute.
Diffstat (limited to 'clang/lib/Frontend/CompilerInvocation.cpp')
0 files changed, 0 insertions, 0 deletions