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