[clang] f3b29c3 - [CIR][AMDGPU] Honor __launch_bounds__ on HIP kernels (#233929)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 21 23:25:06 PDT 2026
Author: Arseniy Obolenskiy
Date: 2026-10-22T08:25:00+02:00
New Revision: f3b29c37148ae577fca3850bb8e3a18bd96d69a1
URL: https://github.com/llvm/llvm-project/commit/f3b29c37148ae577fca3850bb8e3a18bd96d69a1
DIFF: https://github.com/llvm/llvm-project/commit/f3b29c37148ae577fca3850bb8e3a18bd96d69a1.diff
LOG: [CIR][AMDGPU] Honor __launch_bounds__ on HIP kernels (#233929)
Map max threads to amdgpu-flat-work-group-size and min blocks to
amdgpu-waves-per-eu matching classic codegen
Added:
Modified:
clang/include/clang/CIR/Dialect/IR/CIRDialect.td
clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp
clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
Removed:
################################################################################
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
index 53e4730e95d978..cd36d618ca14b3 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
@@ -107,6 +107,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 2a4b61b8dc8314..bc5f6b5d0d4569 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;
@@ -58,11 +59,14 @@ 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::CIRDialect::getAMDGPUFlatWorkGroupSizeAttrName(),
+ 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.
@@ -72,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));
}
}
@@ -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::CIRDialect::getAMDGPUWavesPerEUAttrName(),
+ builder.getStringAttr(llvm::utostr(lbMinWaves)));
return;
+ }
unsigned min =
attr->getMin()->EvaluateKnownConstInt(cgm.getASTContext()).getExtValue();
unsigned max = attr->getMax()
@@ -97,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");
}
@@ -226,9 +237,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
More information about the cfe-commits
mailing list