Author: David Rivera Date: 2026-09-04T11:51:52-04:00 New Revision: 9cda50c8427fc17663f102392729bbf68cd8178e
URL: https://github.com/llvm/llvm-project/commit/9cda50c8427fc17663f102392729bbf68cd8178e DIFF: https://github.com/llvm/llvm-project/commit/9cda50c8427fc17663f102392729bbf68cd8178e.diff LOG: [CIR] Implement "uniform_work_group_size" func attribute. (#214708) Added: clang/test/CIR/CodeGenCUDA/uniform-work-group-size.cu Modified: clang/include/clang/CIR/Dialect/IR/CIRDialect.td clang/lib/CIR/CodeGen/CIRGenCall.cpp clang/test/CIR/CodeGenCUDA/kernel-call.cu mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp mlir/lib/Target/LLVMIR/ModuleImport.cpp mlir/lib/Target/LLVMIR/ModuleTranslation.cpp mlir/test/Dialect/LLVMIR/func.mlir mlir/test/Dialect/LLVMIR/roundtrip.mlir mlir/test/Target/LLVMIR/Import/function-attributes.ll mlir/test/Target/LLVMIR/Import/instructions.ll mlir/test/Target/LLVMIR/llvmir.mlir Removed: ################################################################################ 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..752d4d863cb7e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp @@ -433,8 +433,12 @@ 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. + + // -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()); 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..86e47e52dfc7c --- /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(); + +__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 + +// 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]+]] +// 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..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, @@ -877,6 +878,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 +2137,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..cd6425f4a63fa 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, @@ -1603,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, @@ -1614,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, @@ -1624,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 aa62d0d0db4b1..497b24c94811c 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", @@ -665,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 b6ef1503acc22..bc4ca805b9d7b 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)) @@ -3237,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); } @@ -3256,6 +3262,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..42581ca2e3d92 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"} : () -> () @@ -664,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/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..4116bed6f3489 100644 --- a/mlir/test/Target/LLVMIR/Import/instructions.ll +++ b/mlir/test/Target/LLVMIR/Import/instructions.ll @@ -886,6 +886,36 @@ 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() +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() + ; 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..9f99bf83dae95 100644 --- a/mlir/test/Target/LLVMIR/llvmir.mlir +++ b/mlir/test/Target/LLVMIR/llvmir.mlir @@ -3051,6 +3051,52 @@ 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() +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 // CHECK-SAME: #[[ATTRS:[0-9]+]] llvm.func @zero_call_used_regs_1() attributes { zero_call_used_regs = "skip"} { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
