[all-commits] [llvm/llvm-project] 66acb2: [clang][CodeGen][SPIRV] Translate `amdgpu_flat_wor...

[Alex Voicu via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 66acb2694655321b37a1ee3ff19207a111756562
      https://github.com/llvm/llvm-project/commit/66acb2694655321b37a1ee3ff19207a111756562
  Author: Alex Voicu <alexandru.voicu at amd.com>
  Date:   2025-01-07 (Tue, 07 Jan 2025)

  Changed paths:
    M clang/lib/CodeGen/Targets/SPIR.cpp
    M clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
    M clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu

  Log Message:
  -----------
  [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.



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications