[Mlir-commits] [clang] [mlir] [CIR] Implement "uniform_work_group_size" func attribute. (PR #214708)
David Rivera
llvmlistbot at llvm.org
Fri Sep 4 06:49:22 PDT 2026
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 <davidriverg at gmail.com>
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 <davidriverg at gmail.com>
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
More information about the Mlir-commits
mailing list