https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/233929
Map max threads to amdgpu-flat-work-group-size and min blocks to amdgpu-waves-per-eu matching classic codegen >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] [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 2a4b61b8dc831..af9e229c2c653 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 34d4c57874503..9f50a15ba7253 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
