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 <[email protected]> Date: Tue, 11 Aug 2026 18:00:12 +0200 Subject: [PATCH] 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) <[email protected]> --- 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); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
