[clang] [CIR][AMDGPU] Honor __launch_bounds__ on HIP kernels (PR #233929)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 21 00:05:43 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/233929
>From c0c27ec9478f69319d5c8306cb9b848eeefb2136 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 20 Oct 2026 19:50:31 +0200
Subject: [PATCH 1/2] [CIR][AMDGPU] Honor __launch_bounds__ on HIP kernels
Map max threads to amdgpu-flat-work-group-size and min blocks to amdgpu-waves-per-eu matching classic codegen
---
clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp | 33 ++++++++++++++++++----
clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip | 21 ++++++++++++++
2 files changed, 49 insertions(+), 5 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
index 2a4b61b8dc8314..af9e229c2c653b 100644
--- a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
@@ -29,7 +29,8 @@ namespace {
static void
handleAMDGPUFlatWorkGroupSizeAttr(const FunctionDecl *fd, cir::FuncOp func,
CIRGenModule &cgm, CIRGenBuilderTy &builder,
- bool isOpenCLKernel, bool isHIPKernel) {
+ bool isOpenCLKernel, bool isHIPKernel,
+ unsigned lbMaxThreads) {
const auto *flatWGS = fd->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
const auto *reqdWGS =
cgm.getLangOpts().OpenCL ? fd->getAttr<ReqdWorkGroupSizeAttr>() : nullptr;
@@ -63,6 +64,9 @@ handleAMDGPUFlatWorkGroupSizeAttr(const FunctionDecl *fd, cir::FuncOp func,
} else {
assert(max == 0 && "Max must be zero");
}
+ } else if (lbMaxThreads > 0) {
+ func->setAttr("cir.amdgpu-flat-work-group-size",
+ builder.getStringAttr("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.
@@ -80,10 +84,16 @@ handleAMDGPUFlatWorkGroupSizeAttr(const FunctionDecl *fd, cir::FuncOp func,
/// Handle amdgpu-waves-per-eu attribute.
static void handleAMDGPUWavesPerEUAttr(const FunctionDecl *fd, cir::FuncOp func,
CIRGenModule &cgm,
- CIRGenBuilderTy &builder) {
+ CIRGenBuilderTy &builder,
+ unsigned lbMinWaves) {
const auto *attr = fd->getAttr<AMDGPUWavesPerEUAttr>();
- if (!attr)
+ if (!attr) {
+ // HIP maps __launch_bounds__ min blocks to min waves per EU.
+ if (lbMinWaves > 0)
+ func->setAttr("cir.amdgpu-waves-per-eu",
+ builder.getStringAttr(llvm::utostr(lbMinWaves)));
return;
+ }
unsigned min =
attr->getMin()->EvaluateKnownConstInt(cgm.getASTContext()).getExtValue();
unsigned max = attr->getMax()
@@ -226,9 +236,22 @@ void clang::CIRGen::setAMDGPUTargetFunctionAttributes(const Decl *decl,
const bool isHIPKernel =
cgm.getLangOpts().HIP && fd->hasAttr<CUDAGlobalAttr>();
+ // Native AMDGPU attributes take precedence over __launch_bounds__.
+ unsigned lbMaxThreads = 0;
+ unsigned lbMinWaves = 0;
+ if (const auto *lb =
+ isHIPKernel ? fd->getAttr<CUDALaunchBoundsAttr>() : nullptr) {
+ lbMaxThreads = lb->getMaxThreads()
+ ->EvaluateKnownConstInt(cgm.getASTContext())
+ .getExtValue();
+ if (const Expr *minBlocks = lb->getMinBlocks())
+ lbMinWaves =
+ minBlocks->EvaluateKnownConstInt(cgm.getASTContext()).getExtValue();
+ }
+
handleAMDGPUFlatWorkGroupSizeAttr(fd, func, cgm, builder, isOpenCLKernel,
- isHIPKernel);
- handleAMDGPUWavesPerEUAttr(fd, func, cgm, builder);
+ isHIPKernel, lbMaxThreads);
+ handleAMDGPUWavesPerEUAttr(fd, func, cgm, builder, lbMinWaves);
handleAMDGPUNumSGPRAttr(fd, func, cgm, builder);
handleAMDGPUNumVGPRAttr(fd, func, cgm, builder);
handleAMDGPUMaxNumWorkGroupsAttr(fd, func, cgm, builder);
diff --git a/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip b/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
index 34d4c57874503f..9f50a15ba72536 100644
--- a/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
@@ -74,6 +74,24 @@ __attribute__((amdgpu_num_sgpr(48)))
__attribute__((amdgpu_num_vgpr(32)))
__global__ void kernel_combined() {}
+// CIR: cir.func{{.*}} @_Z16kernel_lb_threadv(){{.*}}"cir.amdgpu-flat-work-group-size" = "1,256"
+// CIR-NOT: cir.amdgpu-waves-per-eu
+// LLVM: define{{.*}} void @_Z16kernel_lb_threadv(){{.*}} #[[LB_THREAD_ATTR:[0-9]+]]
+__launch_bounds__(256)
+__global__ void kernel_lb_thread() {}
+
+// CIR: cir.func{{.*}} @_Z15kernel_lb_wavesv(){{.*}}"cir.amdgpu-flat-work-group-size" = "1,128"{{.*}}"cir.amdgpu-waves-per-eu" = "4"
+// LLVM: define{{.*}} void @_Z15kernel_lb_wavesv(){{.*}} #[[LB_WAVES_ATTR:[0-9]+]]
+__launch_bounds__(128, 4)
+__global__ void kernel_lb_waves() {}
+
+// CIR: cir.func{{.*}} @_Z18kernel_lb_overridev(){{.*}}"cir.amdgpu-flat-work-group-size" = "32,64"{{.*}}"cir.amdgpu-waves-per-eu" = "3"
+// LLVM: define{{.*}} void @_Z18kernel_lb_overridev(){{.*}} #[[LB_OVERRIDE_ATTR:[0-9]+]]
+__attribute__((amdgpu_flat_work_group_size(32, 64)))
+__attribute__((amdgpu_waves_per_eu(3)))
+__launch_bounds__(128, 4)
+__global__ void kernel_lb_override() {}
+
// Test: Device function should NOT have kernel attributes
// CIR: cir.func{{.*}} @_Z9device_fnv()
// CIR-NOT: cir.amdgpu-flat-work-group-size
@@ -88,6 +106,9 @@ __device__ void device_fn() {}
// LLVM-DAG: attributes #[[SGPR_ATTR]] = {{.*}}"amdgpu-num-sgpr"="32"
// LLVM-DAG: attributes #[[VGPR_ATTR]] = {{.*}}"amdgpu-num-vgpr"="64"
// LLVM-DAG: attributes #[[MAX_WG_ATTR]] = {{.*}}"amdgpu-max-num-workgroups"="8,4,2"
+// LLVM-DAG: attributes #[[LB_THREAD_ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="1,256"
+// LLVM-DAG: attributes #[[LB_WAVES_ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="1,128"{{.*}}"amdgpu-waves-per-eu"="4"
+// LLVM-DAG: attributes #[[LB_OVERRIDE_ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="32,64"{{.*}}"amdgpu-waves-per-eu"="3"
// LLVM-DAG: attributes #[[COMBINED_ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="256,256"{{.*}}"amdgpu-num-sgpr"="48"{{.*}}"amdgpu-num-vgpr"="32"{{.*}}"amdgpu-waves-per-eu"="1,2"
// Test: amdgpu-expand-waitcnt-profiling is set on all functions when enabled
>From 68ce563e95ae70c09cdf1e4dfbbe46ff00c8a3a3 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 21 Oct 2026 09:05:16 +0200
Subject: [PATCH 2/2] strings
---
clang/include/clang/CIR/Dialect/IR/CIRDialect.td | 2 ++
clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp | 11 ++++++-----
2 files changed, 8 insertions(+), 5 deletions(-)
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
index 5bd26c6fa9b70d..3c70f67038ee4f 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
@@ -105,6 +105,8 @@ def CIR_Dialect : Dialect {
static llvm::StringRef getTargetCPUAttrName() { return "cir.target-cpu"; }
static llvm::StringRef getTuneCPUAttrName() { return "cir.tune-cpu"; }
static llvm::StringRef getTargetFeaturesAttrName() { return "cir.target-features"; }
+ static llvm::StringRef getAMDGPUFlatWorkGroupSizeAttrName() { return "cir.amdgpu-flat-work-group-size"; }
+ static llvm::StringRef getAMDGPUWavesPerEUAttrName() { return "cir.amdgpu-waves-per-eu"; }
// Raw bytes of the device-side fat binary, read by CIRGen so LoweringPrepare
// can build the runtime-registration globals without doing file I/O.
static llvm::StringRef getCUDADeviceBinaryAttrName() { return "cir.cu.device_binary"; }
diff --git a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
index af9e229c2c653b..bc5f6b5d0d4569 100644
--- a/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
@@ -59,13 +59,13 @@ handleAMDGPUFlatWorkGroupSizeAttr(const FunctionDecl *fd, cir::FuncOp func,
if (min != 0) {
assert(min <= max && "Min must be less than or equal Max");
std::string attrVal = llvm::utostr(min) + "," + llvm::utostr(max);
- func->setAttr("cir.amdgpu-flat-work-group-size",
+ func->setAttr(cir::CIRDialect::getAMDGPUFlatWorkGroupSizeAttrName(),
builder.getStringAttr(attrVal));
} else {
assert(max == 0 && "Max must be zero");
}
} else if (lbMaxThreads > 0) {
- func->setAttr("cir.amdgpu-flat-work-group-size",
+ func->setAttr(cir::CIRDialect::getAMDGPUFlatWorkGroupSizeAttrName(),
builder.getStringAttr("1," + llvm::utostr(lbMaxThreads)));
} else if (isOpenCLKernel || isHIPKernel) {
// By default, restrict the maximum size to a value specified by
@@ -76,7 +76,7 @@ handleAMDGPUFlatWorkGroupSizeAttr(const FunctionDecl *fd, cir::FuncOp func,
: cgm.getLangOpts().GPUMaxThreadsPerBlock;
std::string attrVal =
std::string("1,") + llvm::utostr(defaultMaxWorkGroupSize);
- func->setAttr("cir.amdgpu-flat-work-group-size",
+ func->setAttr(cir::CIRDialect::getAMDGPUFlatWorkGroupSizeAttrName(),
builder.getStringAttr(attrVal));
}
}
@@ -90,7 +90,7 @@ static void handleAMDGPUWavesPerEUAttr(const FunctionDecl *fd, cir::FuncOp func,
if (!attr) {
// HIP maps __launch_bounds__ min blocks to min waves per EU.
if (lbMinWaves > 0)
- func->setAttr("cir.amdgpu-waves-per-eu",
+ func->setAttr(cir::CIRDialect::getAMDGPUWavesPerEUAttrName(),
builder.getStringAttr(llvm::utostr(lbMinWaves)));
return;
}
@@ -107,7 +107,8 @@ static void handleAMDGPUWavesPerEUAttr(const FunctionDecl *fd, cir::FuncOp func,
std::string attrVal = llvm::utostr(min);
if (max != 0)
attrVal = attrVal + "," + llvm::utostr(max);
- func->setAttr("cir.amdgpu-waves-per-eu", builder.getStringAttr(attrVal));
+ func->setAttr(cir::CIRDialect::getAMDGPUWavesPerEUAttrName(),
+ builder.getStringAttr(attrVal));
} else {
assert(max == 0 && "Max must be zero");
}
More information about the cfe-commits
mailing list