[Mlir-commits] [mlir] 9cda50c - [CIR] Implement "uniform_work_group_size" func attribute. (#214708)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Sep 4 08:51:57 PDT 2026
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"} {
More information about the Mlir-commits
mailing list