[clang] clang/AMDGPU: Respect __launch_bounds__ attribute (PR #215615)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Thu Aug 13 03:20:43 PDT 2026
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/215615
>From 0b8da7e78ace7e9c7db5b406256f7462aab280e4 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Tue, 11 Aug 2026 18:00:12 +0200
Subject: [PATCH 1/2] clang/AMDGPU: Respect __launch_bounds__ attribute
Currently the HIP headers manually implement this with a
macro setting amdgpu attributes, and the proper clang attribute
is silently ignored. Directly map the proper attribute into
the target IR attributes. The first argument sets
"amdgpu-flat-work-group-size" and the second (reinterpreted by HIP
as minimum waves per EU) sets "amdgpu-waves-per-eu". An explicit
amdgpu_flat_work_group_size / amdgpu_waves_per_eu attribute takes
precedence. This matches the launch_bounds macro in the HIP headers,
which can now be dropped.
The 3rd maxclusterrank argument is only handled for NVPTX, so restrict
the sm_90 arch check to NVPTX targets and ignore the third argument on
other targets.
Fixes #91468
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
---
clang/include/clang/Basic/Attr.td | 5 +-
clang/include/clang/Basic/AttrDocs.td | 43 ++++++++++++
clang/lib/CodeGen/Targets/AMDGPU.cpp | 31 ++++++++-
clang/lib/Sema/SemaDeclAttr.cpp | 23 +++++--
clang/test/AST/undocumented-attrs.cpp | 3 +-
clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu | 68 +++++++++++++++++++
clang/test/SemaCUDA/launch_bounds_amdgpu.cu | 16 +++++
7 files changed, 176 insertions(+), 13 deletions(-)
create mode 100644 clang/test/SemaCUDA/launch_bounds_amdgpu.cu
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 7d66e42700eef..f599e6fecc8ff 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1653,10 +1653,7 @@ def CUDALaunchBounds : InheritableAttr {
ExprArgument<"MaxBlocks", 1>];
let LangOpts = [CUDA];
let Subjects = SubjectList<[ObjCMethod, FunctionLike]>;
- // An AST node is created for this attribute, but is not used by other parts
- // of the compiler. However, this node needs to exist in the AST because
- // non-LLVM backends may be relying on the attribute's presence.
- let Documentation = [Undocumented];
+ let Documentation = [LaunchBoundsDocs];
}
def CUDAShared : InheritableAttr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 847e44688ed95..7e2ad1754a850 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -3554,6 +3554,49 @@ create machine code that meets the request.
}];
}
+def LaunchBoundsDocs : Documentation {
+ let Category = DocCatFunction;
+ let Content = [{
+The ``__launch_bounds__`` attribute (also spelled ``launch_bounds``) originates
+in CUDA. It informs the compiler of the launch configuration a kernel will be
+dispatched with, allowing it to optimize the kernel accordingly. It takes the
+form ``__launch_bounds__(<max-threads-per-block>[,
+<min-blocks-per-multiprocessor>[, <max-blocks-per-cluster>]])``. All arguments
+are constant expressions.
+
+The attribute only takes effect on ``__global__`` (kernel) functions; like
+NVCC, Clang ignores it on any other function.
+
+``<max-threads-per-block>`` specifies the maximum number of threads per block
+the kernel will be launched with. ``<min-blocks-per-multiprocessor>`` specifies
+the desired minimum number of blocks resident per multiprocessor, and
+``<max-blocks-per-cluster>`` the maximum number of blocks per cluster.
+
+For the NVPTX target, ``<max-threads-per-block>`` and
+``<min-blocks-per-multiprocessor>`` map to the ``.maxntid`` and ``.minnctapersm``
+PTX directives, respectively, and ``<max-blocks-per-cluster>`` (which requires
+``sm_90`` or newer) maps to ``.maxclusterrank``.
+
+For the AMDGPU target, the attribute is translated into the equivalent AMDGPU
+kernel attributes:
+
+ - ``<max-threads-per-block>`` sets the maximum
+ ``amdgpu_flat_work_group_size`` (as ``1, <max-threads-per-block>``).
+ - ``<min-blocks-per-multiprocessor>`` sets the minimum
+ ``amdgpu_waves_per_eu``. Note that HIP reinterprets this CUDA argument as a
+ minimum number of waves per execution unit, so its meaning differs from the
+ NVPTX interpretation.
+ - ``<max-blocks-per-cluster>`` is currently ignored.
+
+An explicit ``amdgpu_flat_work_group_size`` or ``amdgpu_waves_per_eu`` attribute
+takes precedence over the value derived from ``__launch_bounds__``.
+
+When the same kernel is declared multiple times, the launch bounds from the most
+recent declaration that specifies them are used; a definition without
+``__launch_bounds__`` inherits the bounds from an earlier declaration.
+ }];
+}
+
def AMDGPUNumSGPRNumVGPRDocs : Documentation {
let Category = DocCatAMDGPUAttributes;
let Content = [{
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index 9a8fcc546915e..9068c42287c0a 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -347,8 +347,30 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
const bool IsHIPKernel = M.getLangOpts().HIP && FD->hasAttr<CUDAGlobalAttr>();
const auto *FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
+
+ // __launch_bounds__ only takes effect on kernels and is silently ignored on
+ // othre functions The arguments are honored only if the equivalent native
+ // amdgpu_flat_work_group_size / amdgpu_waves_per_eu attribute was not also
+ // used out; those take precedence.
+ const auto *LaunchBounds =
+ IsHIPKernel ? FD->getAttr<CUDALaunchBoundsAttr>() : nullptr;
+ unsigned LBMaxThreads = 0;
+ unsigned LBMinWaves = 0;
+ if (LaunchBounds) {
+ LBMaxThreads = LaunchBounds->getMaxThreads()
+ ->EvaluateKnownConstInt(M.getContext())
+ .getExtValue();
+ if (const Expr *MinBlocks = LaunchBounds->getMinBlocks()) {
+ LBMinWaves =
+ MinBlocks->EvaluateKnownConstInt(M.getContext()).getExtValue();
+ }
+ }
+
if (ReqdWGS || FlatWGS) {
M.handleAMDGPUFlatWorkGroupSizeAttr(F, FlatWGS, ReqdWGS);
+ } else if (LBMaxThreads > 0) {
+ F->addFnAttr("amdgpu-flat-work-group-size",
+ "1," + llvm::utostr(LBMaxThreads));
} else if (IsOpenCLKernel || IsHIPKernel) {
// By default, restrict the maximum size to a value specified by
// --gpu-max-threads-per-block=n or its default value for HIP.
@@ -361,8 +383,15 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
F->addFnAttr("amdgpu-flat-work-group-size", AttrVal);
}
- if (const auto *Attr = FD->getAttr<AMDGPUWavesPerEUAttr>())
+ if (const auto *Attr = FD->getAttr<AMDGPUWavesPerEUAttr>()) {
M.handleAMDGPUWavesPerEUAttr(F, Attr);
+ } else if (LBMinWaves > 0) {
+ // HIP reinterprets the second argument as the minimum waves per EU.
+ //
+ // TODO: The third argument (maxclusterrank) could be used if the AMDGPU
+ // "clusters" feature is supported for the current subtarget.
+ F->addFnAttr("amdgpu-waves-per-eu", llvm::utostr(LBMinWaves));
+ }
if (const auto *Attr = FD->getAttr<AMDGPUNumSGPRAttr>()) {
unsigned NumSGPR = Attr->getNumSGPR();
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 0645f99492433..d19a2e531a9a3 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -6128,12 +6128,23 @@ Sema::CreateLaunchBoundsAttr(const AttributeCommonInfo &CI, Expr *MaxThreads,
// We might want to ignore the nvptx arch check, e.g., when processing the
// launch bounds attribute within ompx_attribute to support other archs.
if (!IgnoreArch) {
- // '.maxclusterrank' ptx directive requires .target sm_90 or higher.
- OffloadArch SM = getOffloadArch(Context.getTargetInfo());
- if (SM.isUnknown() || llvm::NVPTX::getSmVersion(SM.nvptxKind()) < 900) {
- Diag(MaxBlocks->getBeginLoc(), diag::warn_cuda_maxclusterrank_sm_90)
- << OffloadArchToString(SM) << CI << MaxBlocks->getSourceRange();
- // Ignore it by setting MaxBlocks to null;
+ const TargetInfo &DeviceTI =
+ (!Context.getLangOpts().CUDAIsDevice && Context.getAuxTargetInfo())
+ ? *Context.getAuxTargetInfo()
+ : Context.getTargetInfo();
+ if (DeviceTI.getTriple().isNVPTX()) {
+ // '.maxclusterrank' ptx directive requires .target sm_90 or higher.
+ OffloadArch SM = getOffloadArch(DeviceTI);
+ if (SM.isUnknown() || llvm::NVPTX::getSmVersion(SM.nvptxKind()) < 900) {
+ Diag(MaxBlocks->getBeginLoc(), diag::warn_cuda_maxclusterrank_sm_90)
+ << OffloadArchToString(SM) << CI << MaxBlocks->getSourceRange();
+ // Ignore it by setting MaxBlocks to null;
+ MaxBlocks = nullptr;
+ }
+ } else {
+ // maxclusterrank is only handled for NVPTX; ignore it elsewhere.
+ // TODO: Interpret this for AMDGPU with the "clusters" subtarget
+ // feature.
MaxBlocks = nullptr;
}
}
diff --git a/clang/test/AST/undocumented-attrs.cpp b/clang/test/AST/undocumented-attrs.cpp
index 64c7d165915cc..525466b1efdc0 100644
--- a/clang/test/AST/undocumented-attrs.cpp
+++ b/clang/test/AST/undocumented-attrs.cpp
@@ -21,7 +21,6 @@ CHECK-NEXT: CUDAConstant
CHECK-NEXT: CUDADevice
CHECK-NEXT: CUDAGlobal
CHECK-NEXT: CUDAHost
-CHECK-NEXT: CUDALaunchBounds
CHECK-NEXT: CUDAShared
CHECK-NEXT: Capability
CHECK-NEXT: Common
@@ -89,4 +88,4 @@ CHECK-NEXT: Visibility
CHECK-NEXT: WeakImport
CHECK-NEXT: WeakRef
CHECK-NEXT: WorkGroupSizeHint
-CHECK-NEXT: Total: 83
+CHECK-NEXT: Total: 82
diff --git a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
index c1ba7a178cea5..da88c8bd6ebc7 100644
--- a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
@@ -96,6 +96,66 @@ __global__ void template_a_b_c_max_num_work_groups() {}
template __global__ void template_a_b_c_max_num_work_groups<32, 4, 2>();
// CHECK: define{{.*}} amdgpu_kernel void @_Z34template_a_b_c_max_num_work_groupsILj32ELj4ELj2EEvv() [[MAX_NUM_WORK_GROUPS_32_4_2]]
+// __launch_bounds__ is consumed directly on AMDGPU: the first argument maps to
+// the maximum flat work group size and the (optional) second to the minimum
+// waves per execution unit.
+__launch_bounds__(128)
+__global__ void launch_bounds_1arg() {
+// CHECK: define{{.*}} amdgpu_kernel void @_Z18launch_bounds_1argv() [[LAUNCH_BOUNDS_1ARG:#[0-9]+]]
+}
+
+__launch_bounds__(128, 2)
+__global__ void launch_bounds_2arg() {
+// CHECK: define{{.*}} amdgpu_kernel void @_Z18launch_bounds_2argv() [[LAUNCH_BOUNDS_2ARG:#[0-9]+]]
+}
+
+// The third argument (maxclusterrank) is not yet handled on AMDGPU; it is
+// silently ignored without the NVPTX sm_90 diagnostic.
+__launch_bounds__(128, 2, 4)
+__global__ void launch_bounds_3arg() {
+// CHECK: define{{.*}} amdgpu_kernel void @_Z18launch_bounds_3argv() [[LAUNCH_BOUNDS_2ARG]]
+}
+
+// An explicit amdgpu_flat_work_group_size / amdgpu_waves_per_eu takes precedence
+// over __launch_bounds__.
+__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]+]]
+}
+
+// The launch bounds from an earlier declaration are kept when the definition
+// does not specify any.
+__launch_bounds__(128, 2)
+__global__ void launch_bounds_redecl_def_none();
+__global__ void launch_bounds_redecl_def_none() {
+// CHECK: define{{.*}} amdgpu_kernel void @_Z29launch_bounds_redecl_def_nonev() [[LAUNCH_BOUNDS_2ARG]]
+}
+
+// Launch bounds specified only on the definition are honored.
+__global__ void launch_bounds_redecl_decl_none();
+__launch_bounds__(128, 2)
+__global__ void launch_bounds_redecl_decl_none() {
+// CHECK: define{{.*}} amdgpu_kernel void @_Z30launch_bounds_redecl_decl_nonev() [[LAUNCH_BOUNDS_2ARG]]
+}
+
+// When multiple declarations specify conflicting launch bounds, the last one
+// wins.
+__launch_bounds__(64, 8)
+__global__ void launch_bounds_redecl_conflict();
+__launch_bounds__(128, 2)
+__global__ void launch_bounds_redecl_conflict();
+__global__ void launch_bounds_redecl_conflict() {
+// CHECK: define{{.*}} amdgpu_kernel void @_Z29launch_bounds_redecl_conflictv() [[LAUNCH_BOUNDS_2ARG]]
+}
+
+// __launch_bounds__ only takes effect on kernels; it is silently ignored on
+// __device__ functions.
+__launch_bounds__(128, 2)
+__device__ void launch_bounds_device_fn() {
+// CHECK: define{{.*}} void @_Z23launch_bounds_device_fnv() [[LAUNCH_BOUNDS_DEVICE:#[0-9]+]]
+}
+
// Make sure this is silently accepted on other targets.
// NAMD-NOT: "amdgpu-flat-work-group-size"
// NAMD-NOT: "amdgpu-waves-per-eu"
@@ -113,5 +173,13 @@ template __global__ void template_a_b_c_max_num_work_groups<32, 4, 2>();
// CHECK-DAG: attributes [[NUM_VGPR_64]] = {{.*}}"amdgpu-num-vgpr"="64"
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_4_2]] = {{.*}}"amdgpu-max-num-workgroups"="32,4,2"
// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_1_1]] = {{.*}}"amdgpu-max-num-workgroups"="32,1,1"
+// CHECK-DAG: attributes [[LAUNCH_BOUNDS_1ARG]] = {{.*}}"amdgpu-flat-work-group-size"="1,128"
+// CHECK-DAG: attributes [[LAUNCH_BOUNDS_2ARG]] = {{.*}}"amdgpu-flat-work-group-size"="1,128"{{.*}}"amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[LAUNCH_BOUNDS_OVERRIDE]] = {{.*}}"amdgpu-flat-work-group-size"="8,32"{{.*}}"amdgpu-waves-per-eu"="4"
+// __launch_bounds__ is ignored on __device__ functions, so no
+// amdgpu-flat-work-group-size / amdgpu-waves-per-eu attribute is present.
+// 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" }
// NOUB-NOT: "uniform-work-group-size"
diff --git a/clang/test/SemaCUDA/launch_bounds_amdgpu.cu b/clang/test/SemaCUDA/launch_bounds_amdgpu.cu
new file mode 100644
index 0000000000000..8f7ea10a490b7
--- /dev/null
+++ b/clang/test/SemaCUDA/launch_bounds_amdgpu.cu
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-is-device \
+// RUN: -triple amdgpu9.0a-amd-amdhsa -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-is-device \
+// RUN: -triple spirv64-amd-amdhsa -verify %s
+
+// expected-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+// The one- and two-argument forms are consumed by AMDGPU codegen.
+__launch_bounds__(128) void Test1Arg(void);
+__launch_bounds__(128, 2) void Test2Args(void);
+
+// The third argument (maxclusterrank) is not yet handled on AMDGPU; it is
+// silently ignored rather than triggering the NVPTX sm_90 diagnostic.
+__launch_bounds__(128, 2, 4) void Test3Args(void);
>From 4a2236dffc49d1c272ed8629b6200639c041592e Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 13 Aug 2026 12:20:33 +0200
Subject: [PATCH 2/2] Typo fix
Co-authored-by: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
---
clang/lib/CodeGen/Targets/AMDGPU.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index 9068c42287c0a..07e2eac39305d 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -349,7 +349,7 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
const auto *FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
// __launch_bounds__ only takes effect on kernels and is silently ignored on
- // othre functions The arguments are honored only if the equivalent native
+ // other functions The arguments are honored only if the equivalent native
// amdgpu_flat_work_group_size / amdgpu_waves_per_eu attribute was not also
// used out; those take precedence.
const auto *LaunchBounds =
More information about the cfe-commits
mailing list