https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/214708
>From 67c3d043bcfde8cf7e993cea69705ef841ddc63f Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Fri, 7 Aug 2026 07:30:05 -0400 Subject: [PATCH 1/2] [CIR] Implement "uniform_work_group_size" func attribute. --- .../clang/CIR/Dialect/IR/CIRDialect.td | 1 + clang/lib/CIR/CodeGen/CIRGenCall.cpp | 11 ++- clang/test/CIR/CodeGenCUDA/kernel-call.cu | 16 ++-- .../CodeGenCUDA/uniform-work-group-size.cu | 84 +++++++++++++++++++ mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td | 2 + mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 4 + .../LLVMIR/LLVMToLLVMIRTranslation.cpp | 3 + mlir/lib/Target/LLVMIR/ModuleImport.cpp | 5 ++ mlir/lib/Target/LLVMIR/ModuleTranslation.cpp | 2 + mlir/test/Dialect/LLVMIR/func.mlir | 6 ++ mlir/test/Dialect/LLVMIR/roundtrip.mlir | 3 + .../LLVMIR/Import/function-attributes.ll | 6 ++ .../test/Target/LLVMIR/Import/instructions.ll | 12 +++ mlir/test/Target/LLVMIR/llvmir.mlir | 27 ++++++ 14 files changed, 172 insertions(+), 10 deletions(-) create mode 100644 clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index 135cbdcc7d7be..9ea186489ac65 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td @@ -76,6 +76,7 @@ def CIR_Dialect : Dialect { static llvm::StringRef getZeroCallUsedRegsAttrName() { return "zero_call_used_regs"; } static llvm::StringRef getSaveRegParamsAttrName() { return "save_reg_params"; } static llvm::StringRef getDefaultFuncAttrsAttrName() { return "default_func_attrs"; } + static llvm::StringRef getUniformWorkGroupSizeAttrName() { return "uniform_work_group_size"; } static llvm::StringRef getResAttrsAttrName() { return "res_attrs"; } static llvm::StringRef getArgAttrsAttrName() { return "arg_attrs"; } static llvm::StringRef getRecordLayoutsAttrName() { return "cir.record_layouts"; } diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp index 3a4b7cecf2e08..6197fecfb6ee0 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp @@ -433,8 +433,15 @@ void CIRGenModule::constructAttributeList( } } - // TODO(cir): Quite a few CUDA and OpenCL attributes are added here, like - // uniform-work-group-size. + // TODO(cir): Quite a few CUDA and OpenCL attributes are added here. + + // OpenCL v2.0 Work groups may be whether uniform or not. + // '-cl-uniform-work-group-size' compile option gets a hint + // to the compiler that the global work-size be a multiple of + // the work-group size specified to clEnqueueNDRangeKernel + // (i.e. work groups are uniform). + if (langOpts.OffloadUniformBlock) + addUnitAttr(cir::CIRDialect::getUniformWorkGroupSizeAttrName()); if (langOpts.CUDA && !langOpts.CUDAIsDevice && targetDecl->hasAttr<CUDAGlobalAttr>()) { diff --git a/clang/test/CIR/CodeGenCUDA/kernel-call.cu b/clang/test/CIR/CodeGenCUDA/kernel-call.cu index 333af22473a02..9848a48bedb7b 100644 --- a/clang/test/CIR/CodeGenCUDA/kernel-call.cu +++ b/clang/test/CIR/CodeGenCUDA/kernel-call.cu @@ -94,10 +94,10 @@ int main(void) { // HIP-NEW-DAG: cir.alloca "agg.tmp1" {{.*}} : !cir.ptr<!rec_dim3> // // Check dim3 constructors are called for grid and block dimensions - // CUDA-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () - // CUDA-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () - // HIP-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () - // HIP-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () + // CUDA-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) {uniform_work_group_size} : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () + // CUDA-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) {uniform_work_group_size} : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () + // HIP-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) {uniform_work_group_size} : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () + // HIP-NEW: cir.call @_ZN4dim3C1Ejjj({{.*}}) {uniform_work_group_size} : (!cir.ptr<!rec_dim3> {llvm.align = 4 : i64, llvm.dereferenceable = 12 : i64, llvm.nonnull, llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}, !u32i {llvm.noundef}) -> () // // Check default shared memory (0) and null stream are set // CUDA-NEW: cir.const #cir.int<0> : !u64i @@ -106,8 +106,8 @@ int main(void) { // HIP-NEW: cir.const #cir.ptr<null> : !cir.ptr<!rec_hipStream> // // Check Push call configuration is called with grid, block, shared mem, stream - // CUDA-NEW: cir.call @__cudaPushCallConfiguration({{.*}}) : (!u64i, !u32i, !u64i, !u32i, !u64i {llvm.noundef}, !cir.ptr<!rec_cudaStream> {llvm.noundef}) -> !s32i - // HIP-NEW: cir.call @__hipPushCallConfiguration({{.*}}) : (!u64i, !u32i, !u64i, !u32i, !u64i {llvm.noundef}, !cir.ptr<!rec_hipStream> {llvm.noundef}) -> !u32i + // CUDA-NEW: cir.call @__cudaPushCallConfiguration({{.*}}) {uniform_work_group_size} : (!u64i, !u32i, !u64i, !u32i, !u64i {llvm.noundef}, !cir.ptr<!rec_cudaStream> {llvm.noundef}) -> !s32i + // HIP-NEW: cir.call @__hipPushCallConfiguration({{.*}}) {uniform_work_group_size} : (!u64i, !u32i, !u64i, !u32i, !u64i {llvm.noundef}, !cir.ptr<!rec_hipStream> {llvm.noundef}) -> !u32i // // Check the config result is cast to bool for the conditional // CUDA-NEW: cir.cast int_to_bool {{.*}} : !s32i -> !cir.bool @@ -118,13 +118,13 @@ int main(void) { // CUDA-NEW: } else { // CUDA-NEW: cir.const #cir.int<42> : !s32i // CUDA-NEW: cir.const #cir.fp<1.000000e+00> : !cir.float - // CUDA-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {cu.kernel_name = #cir.cu.kernel_name<"_Z6kernelif">} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> () + // CUDA-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {cu.kernel_name = #cir.cu.kernel_name<"_Z6kernelif">, uniform_work_group_size} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> () // CUDA-NEW: } // HIP-NEW: cir.if %{{.*}} { // HIP-NEW: } else { // HIP-NEW: cir.const #cir.int<42> : !s32i // HIP-NEW: cir.const #cir.fp<1.000000e+00> : !cir.float - // HIP-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {cu.kernel_name = #cir.cu.kernel_name<"_Z6kernelif">} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> () + // HIP-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {cu.kernel_name = #cir.cu.kernel_name<"_Z6kernelif">, uniform_work_group_size} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> () // HIP-NEW: } kernel<<<1, 1>>>(42, 1.0f); } diff --git a/clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu b/clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu new file mode 100644 index 0000000000000..f0ae45de48c6d --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu @@ -0,0 +1,84 @@ +// Based on the 'uniform-work-group-size' portion of +// clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu and +// clang/test/CodeGenHIP/default-attributes.hip + +// 'uniform-work-group-size' comes from a language option rather than a decl +// attribute, so it lands on every function and every call site. It defaults to +// on for CUDA/HIP. + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s --check-prefix=CIR +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -foffload-uniform-block -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s --check-prefix=CIR + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s --check-prefix=LLVM +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -foffload-uniform-block -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s --check-prefix=LLVM + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s --check-prefix=OGCG +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -foffload-uniform-block -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s --check-prefix=OGCG + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -fno-offload-uniform-block -fclangir -emit-cir %s -o %t-noub.cir +// RUN: FileCheck --input-file=%t-noub.cir %s --check-prefix=CIR-NOUB +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -fno-offload-uniform-block -fclangir -emit-llvm %s -o %t-noub-cir.ll +// RUN: FileCheck --input-file=%t-noub-cir.ll %s --check-prefix=NOUB +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -fno-offload-uniform-block -emit-llvm %s -o %t-noub.ll +// RUN: FileCheck --input-file=%t-noub.ll %s --check-prefix=NOUB + +#include "Inputs/cuda.h" + +__device__ void extern_func(); + +// CIR: cir.func private @_Z11extern_funcv() +// CIR-SAME: uniform_work_group_size + +__device__ void func() { + extern_func(); +} +// CIR: cir.func{{.*}}@_Z4funcv() +// CIR-SAME: uniform_work_group_size +// CIR: cir.call @_Z11extern_funcv() +// CIR-SAME: uniform_work_group_size + +// LLVM: define{{.*}} void @_Z4funcv() [[FUNC:#[0-9]+]] +// LLVM: call void @_Z11extern_funcv() [[CALL:#[0-9]+]] +// OGCG: define{{.*}} void @_Z4funcv() [[FUNC:#[0-9]+]] +// OGCG: call void @_Z11extern_funcv() [[CALL:#[0-9]+]] + +__global__ void kernel() { + extern_func(); +} +// CIR: cir.func{{.*}}@_Z6kernelv() cc(amdgpu_kernel) +// CIR-SAME: uniform_work_group_size +// CIR: cir.call @_Z11extern_funcv() +// CIR-SAME: uniform_work_group_size + +// LLVM: define{{.*}} amdgpu_kernel void @_Z6kernelv() [[KERNEL:#[0-9]+]] +// LLVM: call void @_Z11extern_funcv() [[CALL]] +// OGCG: define{{.*}} amdgpu_kernel void @_Z6kernelv() [[KERNEL:#[0-9]+]] +// OGCG: call void @_Z11extern_funcv() [[CALL]] + +// The attribute is present on both definitions and call sites. + +// LLVM-DAG: attributes [[FUNC]] = {{.*}}"uniform-work-group-size" +// LLVM-DAG: attributes [[KERNEL]] = {{.*}}"uniform-work-group-size" +// LLVM-DAG: attributes [[CALL]] = {{.*}}"uniform-work-group-size" + +// OGCG-DAG: attributes [[FUNC]] = {{.*}}"uniform-work-group-size" +// OGCG-DAG: attributes [[KERNEL]] = {{.*}}"uniform-work-group-size" +// OGCG-DAG: attributes [[CALL]] = {{.*}}"uniform-work-group-size" + +// CIR-NOUB-NOT: uniform_work_group_size +// NOUB-NOT: "uniform-work-group-size" diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td index 42b10d3df6d2c..0b4dff380fe27 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td @@ -877,6 +877,7 @@ def LLVM_CallOp OptionalAttr<StrAttr>:$zero_call_used_regs, OptionalAttr<StrAttr>:$trap_func_name, OptionalAttr<DictionaryAttr>:$default_func_attrs, + UnitAttr:$uniform_work_group_size, VariadicOfVariadic<LLVM_Type, "op_bundle_sizes">:$op_bundle_operands, DenseI32ArrayAttr:$op_bundle_sizes, OptionalAttr<ArrayAttr>:$op_bundle_tags, @@ -2135,6 +2136,7 @@ def LLVM_LLVMFuncOp : LLVM_Op<"func", [ OptionalAttr<UnitAttr>:$save_reg_params, OptionalAttr<StrAttr>:$zero_call_used_regs, OptionalAttr<DictionaryAttr>:$default_func_attrs, + OptionalAttr<UnitAttr>:$uniform_work_group_size, OptionalAttr<LLVM_VecTypeHintAttr>:$vec_type_hint, OptionalAttr<DenseI32ArrayAttr>:$work_group_size_hint, OptionalAttr<DenseI32ArrayAttr>:$reqd_work_group_size, diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp index 4b4d0f098e559..6093bf29d7ff9 100644 --- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp +++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp @@ -1012,6 +1012,7 @@ void CallOp::build(OpBuilder &builder, OperationState &state, TypeRange results, /*save_reg_params=*/nullptr, /*zero_call_used_regs=*/nullptr, /*trap_func_name=*/nullptr, /*default_func_attrs=*/nullptr, + /*uniform_work_group_size=*/nullptr, /*op_bundle_operands=*/{}, /*op_bundle_tags=*/{}, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, /*access_groups=*/nullptr, /*alias_scopes=*/nullptr, @@ -1052,6 +1053,7 @@ void CallOp::build(OpBuilder &builder, OperationState &state, /*save_reg_params=*/nullptr, /*zero_call_used_regs=*/nullptr, /*trap_func_name=*/nullptr, /*default_func_attrs=*/nullptr, + /*uniform_work_group_size=*/nullptr, /*op_bundle_operands=*/{}, /*op_bundle_tags=*/{}, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, /*access_groups=*/nullptr, @@ -1078,6 +1080,7 @@ void CallOp::build(OpBuilder &builder, OperationState &state, /*save_reg_params=*/nullptr, /*zero_call_used_regs=*/nullptr, /*trap_func_name=*/nullptr, /*default_func_attrs=*/nullptr, + /*uniform_work_group_size=*/nullptr, /*op_bundle_operands=*/{}, /*op_bundle_tags=*/{}, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, /*access_groups=*/nullptr, /*alias_scopes=*/nullptr, @@ -1104,6 +1107,7 @@ void CallOp::build(OpBuilder &builder, OperationState &state, LLVMFuncOp func, /*save_reg_params=*/nullptr, /*zero_call_used_regs=*/nullptr, /*trap_func_name=*/nullptr, /*default_func_attrs=*/nullptr, + /*uniform_work_group_size=*/nullptr, /*op_bundle_operands=*/{}, /*op_bundle_tags=*/{}, /*access_groups=*/nullptr, /*alias_scopes=*/nullptr, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp index aa62d0d0db4b1..78979929a2583 100644 --- a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp +++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp @@ -511,6 +511,9 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder, call->addFnAttr(llvm::Attribute::get(moduleTranslation.getLLVMContext(), "zero-call-used-regs", zcsr.getValue())); + if (callOp.getUniformWorkGroupSizeAttr()) + call->addFnAttr(llvm::Attribute::get(moduleTranslation.getLLVMContext(), + "uniform-work-group-size")); if (StringAttr trapFunc = callOp.getTrapFuncNameAttr()) call->addFnAttr(llvm::Attribute::get(moduleTranslation.getLLVMContext(), "trap-func-name", diff --git a/mlir/lib/Target/LLVMIR/ModuleImport.cpp b/mlir/lib/Target/LLVMIR/ModuleImport.cpp index b6ef1503acc22..b6928182c3e76 100644 --- a/mlir/lib/Target/LLVMIR/ModuleImport.cpp +++ b/mlir/lib/Target/LLVMIR/ModuleImport.cpp @@ -2925,6 +2925,7 @@ static constexpr std::array kExplicitLLVMFuncOpAttributes{ StringLiteral("target-features"), StringLiteral("trap-func-name"), StringLiteral("tune-cpu"), + StringLiteral("uniform-work-group-size"), StringLiteral("uwtable"), StringLiteral("vscale_range"), StringLiteral("willreturn"), @@ -3023,6 +3024,8 @@ void ModuleImport::processFunctionAttributes(llvm::Function *func, funcOp.setOptsize(true); if (func->hasFnAttribute("save-reg-params")) funcOp.setSaveRegParams(true); + if (func->hasFnAttribute("uniform-work-group-size")) + funcOp.setUniformWorkGroupSize(true); if (func->hasFnAttribute(llvm::Attribute::MinSize)) funcOp.setMinsize(true); if (func->hasFnAttribute(llvm::Attribute::ReturnsTwice)) @@ -3256,6 +3259,8 @@ LogicalResult ModuleImport::convertCallAttributes(llvm::CallInst *inst, op.setOptsize( callAttrs.getFnAttr(llvm::Attribute::OptimizeForSize).isValid()); op.setSaveRegParams(callAttrs.getFnAttr("save-reg-params").isValid()); + op.setUniformWorkGroupSize( + callAttrs.getFnAttr("uniform-work-group-size").isValid()); op.setBuiltin(callAttrs.getFnAttr(llvm::Attribute::Builtin).isValid()); op.setNobuiltin(callAttrs.getFnAttr(llvm::Attribute::NoBuiltin).isValid()); op.setMinsize(callAttrs.getFnAttr(llvm::Attribute::MinSize).isValid()); diff --git a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp index 26a497d483912..6f3c17b5a0891 100644 --- a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp +++ b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp @@ -1907,6 +1907,8 @@ static void convertFunctionAttributes(ModuleTranslation &mod, LLVMFuncOp func, convertUWTableKindToLLVM(uwTableKindAttr.getUwtableKind())); if (StringAttr zcsr = func.getZeroCallUsedRegsAttr()) llvmFunc->addFnAttr("zero-call-used-regs", zcsr.getValue()); + if (func.getUniformWorkGroupSizeAttr()) + llvmFunc->addFnAttr("uniform-work-group-size"); if (ArrayAttr noBuiltins = func.getNobuiltinsAttr()) { if (noBuiltins.empty()) diff --git a/mlir/test/Dialect/LLVMIR/func.mlir b/mlir/test/Dialect/LLVMIR/func.mlir index c32dd51b3740e..f243200e7ca13 100644 --- a/mlir/test/Dialect/LLVMIR/func.mlir +++ b/mlir/test/Dialect/LLVMIR/func.mlir @@ -408,6 +408,12 @@ module { llvm.return } + llvm.func @uniform_work_group_size() attributes { uniform_work_group_size } { + // CHECK: @uniform_work_group_size + // CHECK-SAME: attributes {uniform_work_group_size} + llvm.return + } + llvm.func @zero_call_used_regs() attributes { zero_call_used_regs="used-gpr-arg"} { // CHECK: @zero_call_used_regs // CHECK-SAME: attributes {zero_call_used_regs = "used-gpr-arg"} diff --git a/mlir/test/Dialect/LLVMIR/roundtrip.mlir b/mlir/test/Dialect/LLVMIR/roundtrip.mlir index b27d07ecdb4f5..56ba74f6d25ca 100644 --- a/mlir/test/Dialect/LLVMIR/roundtrip.mlir +++ b/mlir/test/Dialect/LLVMIR/roundtrip.mlir @@ -176,6 +176,9 @@ func.func @ops(%arg0: i32, %arg1: f32, // CHECK: llvm.call @baz() {save_reg_params} : () -> () llvm.call @baz() {save_reg_params} : () -> () +// CHECK: llvm.call @baz() {uniform_work_group_size} : () -> () + llvm.call @baz() {uniform_work_group_size} : () -> () + // CHECK: llvm.call @baz() {zero_call_used_regs = "all"} : () -> () llvm.call @baz() {zero_call_used_regs="all"} : () -> () diff --git a/mlir/test/Target/LLVMIR/Import/function-attributes.ll b/mlir/test/Target/LLVMIR/Import/function-attributes.ll index 1c9c616ce3c5b..060b4c51baff5 100644 --- a/mlir/test/Target/LLVMIR/Import/function-attributes.ll +++ b/mlir/test/Target/LLVMIR/Import/function-attributes.ll @@ -568,6 +568,12 @@ declare void @save_reg_params() "save-reg-params" // ----- +; CHECK-LABEL: @uniform_work_group_size +; CHECK-SAME: attributes {uniform_work_group_size} +declare void @uniform_work_group_size() "uniform-work-group-size" + +; // ----- + ; CHECK-LABEL: @zero_call_used_regs ; CHECK-SAME: attributes {zero_call_used_regs = "skip"} declare void @zero_call_used_regs() "zero-call-used-regs"="skip" diff --git a/mlir/test/Target/LLVMIR/Import/instructions.ll b/mlir/test/Target/LLVMIR/Import/instructions.ll index 6fae74dcf9895..a506cc9e0ab2b 100644 --- a/mlir/test/Target/LLVMIR/Import/instructions.ll +++ b/mlir/test/Target/LLVMIR/Import/instructions.ll @@ -886,6 +886,18 @@ define void @call_save_reg_params() { ; CHECK: llvm.func @f() declare void @f() +; CHECK-LABEL: @call_uniform_work_group_size +define void @call_uniform_work_group_size() { +; CHECK: llvm.call @f() {uniform_work_group_size} + call void @f() "uniform-work-group-size" + ret void +} + +; // ----- + +; CHECK: llvm.func @f() +declare void @f() + ; CHECK-LABEL: @call_zero_call_used_regs define void @call_zero_call_used_regs() { ; CHECK: llvm.call @f() {zero_call_used_regs = "used"} diff --git a/mlir/test/Target/LLVMIR/llvmir.mlir b/mlir/test/Target/LLVMIR/llvmir.mlir index 5edc4a0fce9b4..1583471ccae9d 100644 --- a/mlir/test/Target/LLVMIR/llvmir.mlir +++ b/mlir/test/Target/LLVMIR/llvmir.mlir @@ -3051,6 +3051,33 @@ llvm.func @save_reg_params_call() { llvm.func @f() +// CHECK-LABEL: @uniform_work_group_size +// CHECK-SAME: #[[ATTRS:[0-9]+]] +llvm.func @uniform_work_group_size() attributes { uniform_work_group_size } { + llvm.return +} + +// CHECK: #[[ATTRS]] +// CHECK-SAME: "uniform-work-group-size" + +// ----- + +llvm.func @f() + +// CHECK-LABEL: @uniform_work_group_size_call +// CHECK: call void @f() #[[ATTRS:[0-9]+]] +llvm.func @uniform_work_group_size_call() { + llvm.call @f() {uniform_work_group_size} : () -> () + llvm.return +} + +// CHECK: #[[ATTRS]] +// CHECK-SAME: "uniform-work-group-size" + +// ----- + +llvm.func @f() + // CHECK-LABEL: @zero_call_used_regs_1 // CHECK-SAME: #[[ATTRS:[0-9]+]] llvm.func @zero_call_used_regs_1() attributes { zero_call_used_regs = "skip"} { >From c49face5df7e54ab3d115369aa9a08b91ffe3c08 Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Thu, 3 Sep 2026 03:28:40 -0400 Subject: [PATCH 2/2] [CIR] Add uniform_work_group_size support for llvm.invoke ( and address nits) --- clang/lib/CIR/CodeGen/CIRGenCall.cpp | 7 ++----- .../CodeGenCUDA/uniform-work-group-size.cu | 6 +++--- mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td | 1 + mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 11 ++++++----- .../LLVMIR/LLVMToLLVMIRTranslation.cpp | 3 +++ mlir/lib/Target/LLVMIR/ModuleImport.cpp | 3 +++ mlir/test/Dialect/LLVMIR/roundtrip.mlir | 11 +++++++++++ .../test/Target/LLVMIR/Import/instructions.ll | 18 ++++++++++++++++++ mlir/test/Target/LLVMIR/llvmir.mlir | 19 +++++++++++++++++++ 9 files changed, 66 insertions(+), 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp index 6197fecfb6ee0..752d4d863cb7e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp @@ -435,11 +435,8 @@ void CIRGenModule::constructAttributeList( // TODO(cir): Quite a few CUDA and OpenCL attributes are added here. - // OpenCL v2.0 Work groups may be whether uniform or not. - // '-cl-uniform-work-group-size' compile option gets a hint - // to the compiler that the global work-size be a multiple of - // the work-group size specified to clEnqueueNDRangeKernel - // (i.e. work groups are uniform). + // -cl-uniform-work-group-size / -foffload-uniform-block: work groups are + // uniform (global work-size is a multiple of work-group size). if (langOpts.OffloadUniformBlock) addUnitAttr(cir::CIRDialect::getUniformWorkGroupSizeAttrName()); diff --git a/clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu b/clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu index f0ae45de48c6d..86e47e52dfc7c 100644 --- a/clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu +++ b/clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu @@ -41,9 +41,6 @@ __device__ void extern_func(); -// CIR: cir.func private @_Z11extern_funcv() -// CIR-SAME: uniform_work_group_size - __device__ void func() { extern_func(); } @@ -52,6 +49,9 @@ __device__ void func() { // CIR: cir.call @_Z11extern_funcv() // CIR-SAME: uniform_work_group_size +// CIR: cir.func private @_Z11extern_funcv() +// CIR-SAME: uniform_work_group_size + // LLVM: define{{.*}} void @_Z4funcv() [[FUNC:#[0-9]+]] // LLVM: call void @_Z11extern_funcv() [[CALL:#[0-9]+]] // OGCG: define{{.*}} void @_Z4funcv() [[FUNC:#[0-9]+]] diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td index 0b4dff380fe27..ab641a5da8c22 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td @@ -744,6 +744,7 @@ def LLVM_InvokeOp OptionalAttr<DenseI32ArrayAttr>:$branch_weights, DefaultValuedAttr<CConv, "CConv::C">:$CConv, OptionalAttr<DictionaryAttr>:$default_func_attrs, + UnitAttr:$uniform_work_group_size, VariadicOfVariadic<LLVM_Type, "op_bundle_sizes">:$op_bundle_operands, DenseI32ArrayAttr:$op_bundle_sizes, diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp index 6093bf29d7ff9..cd6425f4a63fa 100644 --- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp +++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp @@ -1607,8 +1607,8 @@ void InvokeOp::build(OpBuilder &builder, OperationState &state, LLVMFuncOp func, build(builder, state, getCallOpResultTypes(calleeType), getCallOpVarCalleeType(calleeType), SymbolRefAttr::get(func), ops, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, normalOps, unwindOps, - nullptr, nullptr, /*default_func_attrs=*/nullptr, {}, {}, normal, - unwind); + nullptr, nullptr, /*default_func_attrs=*/nullptr, + /*uniform_work_group_size=*/nullptr, {}, {}, normal, unwind); } void InvokeOp::build(OpBuilder &builder, OperationState &state, TypeRange tys, @@ -1618,7 +1618,8 @@ void InvokeOp::build(OpBuilder &builder, OperationState &state, TypeRange tys, build(builder, state, tys, /*var_callee_type=*/nullptr, callee, ops, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, normalOps, unwindOps, nullptr, nullptr, - /*default_func_attrs=*/nullptr, {}, {}, normal, unwind); + /*default_func_attrs=*/nullptr, + /*uniform_work_group_size=*/nullptr, {}, {}, normal, unwind); } void InvokeOp::build(OpBuilder &builder, OperationState &state, @@ -1628,8 +1629,8 @@ void InvokeOp::build(OpBuilder &builder, OperationState &state, build(builder, state, getCallOpResultTypes(calleeType), getCallOpVarCalleeType(calleeType), callee, ops, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, normalOps, unwindOps, - nullptr, nullptr, /*default_func_attrs=*/nullptr, {}, {}, normal, - unwind); + nullptr, nullptr, /*default_func_attrs=*/nullptr, + /*uniform_work_group_size=*/nullptr, {}, {}, normal, unwind); } SuccessorOperands InvokeOp::getSuccessorOperands(unsigned index) { diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp index 78979929a2583..497b24c94811c 100644 --- a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp +++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp @@ -668,6 +668,9 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder, operandsRef.drop_front(), opBundles); } result->setCallingConv(convertCConvToLLVM(invOp.getCConv())); + if (invOp.getUniformWorkGroupSizeAttr()) + result->addFnAttr(llvm::Attribute::get(moduleTranslation.getLLVMContext(), + "uniform-work-group-size")); moduleTranslation.convertFunctionAttrCollection( invOp.getDefaultFuncAttrsAttr(), result, ModuleTranslation::convertDefaultFuncAttr); diff --git a/mlir/lib/Target/LLVMIR/ModuleImport.cpp b/mlir/lib/Target/LLVMIR/ModuleImport.cpp index b6928182c3e76..bc4ca805b9d7b 100644 --- a/mlir/lib/Target/LLVMIR/ModuleImport.cpp +++ b/mlir/lib/Target/LLVMIR/ModuleImport.cpp @@ -3240,6 +3240,9 @@ static LogicalResult convertCallBaseAttributes(llvm::CallBase *inst, Op op) { LogicalResult ModuleImport::convertInvokeAttributes(llvm::InvokeInst *inst, InvokeOp op) { + llvm::AttributeList invokeAttrs = inst->getAttributes(); + op.setUniformWorkGroupSize( + invokeAttrs.getFnAttr("uniform-work-group-size").isValid()); return convertCallBaseAttributes(inst, op); } diff --git a/mlir/test/Dialect/LLVMIR/roundtrip.mlir b/mlir/test/Dialect/LLVMIR/roundtrip.mlir index 56ba74f6d25ca..42581ca2e3d92 100644 --- a/mlir/test/Dialect/LLVMIR/roundtrip.mlir +++ b/mlir/test/Dialect/LLVMIR/roundtrip.mlir @@ -667,6 +667,17 @@ llvm.func @invokeLandingpad() -> i32 attributes { personality = @__gxx_personali llvm.return %0 : i32 } +// CHECK-LABEL: @invokeUniformWorkGroupSize +llvm.func @invokeUniformWorkGroupSize() attributes { personality = @__gxx_personality_v0 } { + // CHECK: llvm.invoke @baz() to ^{{.*}} unwind ^{{.*}} {uniform_work_group_size} + llvm.invoke @baz() to ^bb1 unwind ^bb2 {uniform_work_group_size} : () -> () +^bb1: + llvm.return +^bb2: + %0 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)> + llvm.return +} + // CHECK-LABEL: @useFreezeOp func.func @useFreezeOp(%arg0: i32) { // CHECK: = llvm.freeze %[[ARG0:.*]] : i32 diff --git a/mlir/test/Target/LLVMIR/Import/instructions.ll b/mlir/test/Target/LLVMIR/Import/instructions.ll index a506cc9e0ab2b..4116bed6f3489 100644 --- a/mlir/test/Target/LLVMIR/Import/instructions.ll +++ b/mlir/test/Target/LLVMIR/Import/instructions.ll @@ -895,6 +895,24 @@ define void @call_uniform_work_group_size() { ; // ----- +; CHECK: llvm.func @f() +declare void @f() +declare i32 @__gxx_personality_v0(...) + +; CHECK-LABEL: @invoke_uniform_work_group_size +define void @invoke_uniform_work_group_size() personality ptr @__gxx_personality_v0 { +entry: +; CHECK: llvm.invoke @f() to ^bb1 unwind ^bb2 {uniform_work_group_size} + invoke void @f() "uniform-work-group-size" to label %bb1 unwind label %bb2 +bb1: + ret void +bb2: + %0 = landingpad i32 cleanup + unreachable +} + +; // ----- + ; CHECK: llvm.func @f() declare void @f() diff --git a/mlir/test/Target/LLVMIR/llvmir.mlir b/mlir/test/Target/LLVMIR/llvmir.mlir index 1583471ccae9d..9f99bf83dae95 100644 --- a/mlir/test/Target/LLVMIR/llvmir.mlir +++ b/mlir/test/Target/LLVMIR/llvmir.mlir @@ -3076,6 +3076,25 @@ llvm.func @uniform_work_group_size_call() { // ----- +llvm.func @f() +llvm.func @__gxx_personality_v0(...) -> i32 + +// CHECK-LABEL: @uniform_work_group_size_invoke +// CHECK: invoke void @f() #[[ATTRS:[0-9]+]] +llvm.func @uniform_work_group_size_invoke() attributes {personality = @__gxx_personality_v0} { + llvm.invoke @f() to ^bb2 unwind ^bb1 {uniform_work_group_size} : () -> () +^bb1: + %0 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)> + llvm.return +^bb2: + llvm.return +} + +// CHECK: #[[ATTRS]] +// CHECK-SAME: "uniform-work-group-size" + +// ----- + llvm.func @f() // CHECK-LABEL: @zero_call_used_regs_1 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
