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 <[email protected]> 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 <[email protected]> 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"); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
