[llvm-branch-commits] [clang] clang/SPIRV: Respect __launch_bounds__ for AMDHIP case (PR #215616)

[Matt Arsenault via llvm-branch-commits]

https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/215616

>From 653134fabea768f3d8de0a0f6c4870a85e8b987d Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Tue, 11 Aug 2026 18:37:18 +0200
Subject: [PATCH] clang/SPIRV: Respect __launch_bounds__ for AMDHIP case

Follow the somewhat dodgy logic for packing amdgpu_flat_work_group_size
into the X field of max_work_group_size if the value is provided
to __launch_bounds__. The explicit amdgpu_flat_work_group_size takes
precedence, like in the AMDGPU case.

Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
---
 clang/lib/CodeGen/Targets/SPIR.cpp            | 8 +++++++-
 clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu | 6 ++++++
 2 files changed, 13 insertions(+), 1 deletion(-)

diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp
index a260469133304..177d166d5b85b 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -542,8 +542,14 @@ void SPIRVTargetCodeGenInfo::setTargetAttributes(
     return;
 
   unsigned N = M.getLangOpts().GPUMaxThreadsPerBlock;
-  if (auto FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>())
+  if (auto FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>()) {
     N = FlatWGS->getMax()->EvaluateKnownConstInt(M.getContext()).getExtValue();
+  } else if (auto LB = FD->getAttr<CUDALaunchBoundsAttr>()) {
+    if (uint64_t MaxThreads = LB->getMaxThreads()
+                                  ->EvaluateKnownConstInt(M.getContext())
+                                  .getExtValue())
+      N = MaxThreads;
+  }
 
   // We encode the maximum flat WG size in the first component of the 3D
   // max_work_group_size attribute, which will get reverse translated into the
diff --git a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
index da88c8bd6ebc7..4009c17b308a5 100644
--- a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
@@ -102,11 +102,13 @@ template __global__ void template_a_b_c_max_num_work_groups<32, 4, 2>();
 __launch_bounds__(128)
 __global__ void launch_bounds_1arg() {
 // CHECK: define{{.*}} amdgpu_kernel void @_Z18launch_bounds_1argv() [[LAUNCH_BOUNDS_1ARG:#[0-9]+]]
+// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z18launch_bounds_1argv(){{.*}} !max_work_group_size [[MAX_WORK_GROUP_SIZE_128:![0-9]+]]
 }
 
 __launch_bounds__(128, 2)
 __global__ void launch_bounds_2arg() {
 // CHECK: define{{.*}} amdgpu_kernel void @_Z18launch_bounds_2argv() [[LAUNCH_BOUNDS_2ARG:#[0-9]+]]
+// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z18launch_bounds_2argv(){{.*}} !max_work_group_size [[MAX_WORK_GROUP_SIZE_128]]
 }
 
 // The third argument (maxclusterrank) is not yet handled on AMDGPU; it is
@@ -114,6 +116,7 @@ __global__ void launch_bounds_2arg() {
 __launch_bounds__(128, 2, 4)
 __global__ void launch_bounds_3arg() {
 // CHECK: define{{.*}} amdgpu_kernel void @_Z18launch_bounds_3argv() [[LAUNCH_BOUNDS_2ARG]]
+// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z18launch_bounds_3argv(){{.*}} !max_work_group_size [[MAX_WORK_GROUP_SIZE_128]]
 }
 
 // An explicit amdgpu_flat_work_group_size / amdgpu_waves_per_eu takes precedence
@@ -122,6 +125,7 @@ __attribute__((amdgpu_flat_work_group_size(8, 32), amdgpu_waves_per_eu(4)))
 __launch_bounds__(128, 2)
 __global__ void launch_bounds_explicit_override() {
 // CHECK: define{{.*}} amdgpu_kernel void @_Z31launch_bounds_explicit_overridev() [[LAUNCH_BOUNDS_OVERRIDE:#[0-9]+]]
+// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z31launch_bounds_explicit_overridev(){{.*}} !max_work_group_size [[MAX_WORK_GROUP_SIZE_32:![0-9]+]]
 }
 
 // The launch bounds from an earlier declaration are kept when the definition
@@ -181,5 +185,7 @@ __device__ void launch_bounds_device_fn() {
 // String attributes are sorted, so the amdgpu-* attributes would appear
 // immediately after "optnone"; check that "no-trapping-math" follows directly.
 // CHECK-DAG: attributes [[LAUNCH_BOUNDS_DEVICE]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"={{.*}}"uniform-work-group-size" }
+// CHECK-SPIRV-DAG: [[MAX_WORK_GROUP_SIZE_128]] = !{i32 128, i32 1, i32 1}
+// CHECK-SPIRV-DAG: [[MAX_WORK_GROUP_SIZE_32]] = !{i32 32, i32 1, i32 1}
 
 // NOUB-NOT: "uniform-work-group-size"