[llvm-branch-commits] [mlir] [mlir][LLVM] Verify that constant attribute and result types match (PR #218887)
Christian Ulmann via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Aug 26 04:57:04 PDT 2026
https://github.com/Dinistro updated https://github.com/llvm/llvm-project/pull/218887
>From 8d735a07a87745a9606ba24a22f201d5d0964e47 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christian.ulmann at nextsilicon.com>
Date: Tue, 25 Aug 2026 09:02:30 +0200
Subject: [PATCH] [mlir][LLVM] Verify that constant attribute and result types
match
`llvm.mlir.constant` only checked that the kind of the value attribute suits
the kind of the result type, never that the types themselves agree. All of the
following verified:
%0 = llvm.mlir.constant(1 : index) : i64
%1 = llvm.mlir.constant(1 : i8) : i16
%2 = llvm.mlir.constant(dense<1> : vector<4xi32>) : vector<4xi64>
Translation ignores the attribute type and uses the result type, so the
attribute type was effectively decorative for integers, and passes that read it
back could observe a type that has nothing to do with the value.
Require exact type equality for integer attributes and exact element type
equality for integer elements attributes, mirroring the `AllTypesMatch`
constraint `arith.constant` gets from ODS. The op cannot use that trait itself
because `value` is an `AnyAttr` that also holds `StringAttr` and `ArrayAttr`.
The element type check is also run on the scalable vector path, which
previously returned early without checking it at all.
Float attributes keep their existing behaviour: an attribute whose element type
has floating point semantics may still be paired with an integer type of the
same width, which is how builtin float types without an LLVM equivalent are
represented.
Note that `index` cannot appear as a constant type, since it is not LLVM
dialect-compatible, so no separate rule is needed for it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
---
mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp | 2 +-
mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp | 2 +-
.../Conversion/ArithToLLVM/ArithToLLVM.cpp | 15 +++--
mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 53 ++++++++++++++----
.../Conversion/ArithToLLVM/arith-to-llvm.mlir | 5 +-
.../ArithToLLVM/constant-index-bitwidth.mlir | 23 ++++++++
mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir | 6 +-
.../OpenMPToLLVM/convert-to-llvmir.mlir | 4 +-
.../Dialect/LLVMIR/inlining-alias-scopes.mlir | 14 ++---
mlir/test/Dialect/LLVMIR/invalid.mlir | 56 +++++++++++++++++++
mlir/test/Dialect/LLVMIR/roundtrip.mlir | 24 ++++----
mlir/test/Dialect/Vector/canonicalize.mlir | 2 +-
mlir/test/Target/LLVMIR/nvvm/fence.mlir | 2 +-
.../Target/LLVMIR/omptarget-if-nowait.mlir | 2 +-
.../LLVMIR/omptarget-parallel-llvm.mlir | 2 +-
.../LLVMIR/openmp-composite-simd-if.mlir | 2 +-
mlir/test/Target/LLVMIR/openmp-llvm.mlir | 2 +-
.../LLVMIR/openmp-target-launch-device.mlir | 8 +--
.../LLVMIR/openmp-target-launch-host.mlir | 4 +-
mlir/test/Transforms/sccp.mlir | 11 ++--
20 files changed, 180 insertions(+), 59 deletions(-)
diff --git a/mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp b/mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp
index 987dfa1eb9e78..a5192900c9206 100644
--- a/mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp
+++ b/mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp
@@ -165,7 +165,7 @@ class PrintOpLowering : public OpConversionPattern<toy::PrintOp> {
// Get the pointer to the first character in the global string.
Value globalPtr = LLVM::AddressOfOp::create(builder, loc, global);
Value cst0 = LLVM::ConstantOp::create(builder, loc, builder.getI64Type(),
- builder.getIndexAttr(0));
+ builder.getI64IntegerAttr(0));
return LLVM::GEPOp::create(
builder, loc, LLVM::LLVMPointerType::get(builder.getContext()),
global.getType(), globalPtr, ArrayRef<Value>({cst0, cst0}));
diff --git a/mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp b/mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp
index 8b48a8f798beb..cf4732ec96a1d 100644
--- a/mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp
+++ b/mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp
@@ -165,7 +165,7 @@ class PrintOpLowering : public OpConversionPattern<toy::PrintOp> {
// Get the pointer to the first character in the global string.
Value globalPtr = LLVM::AddressOfOp::create(builder, loc, global);
Value cst0 = LLVM::ConstantOp::create(builder, loc, builder.getI64Type(),
- builder.getIndexAttr(0));
+ builder.getI64IntegerAttr(0));
return LLVM::GEPOp::create(
builder, loc, LLVM::LLVMPointerType::get(builder.getContext()),
global.getType(), globalPtr, ArrayRef<Value>({cst0, cst0}));
diff --git a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
index 85cf4275e41a8..d164f0e7f6d17 100644
--- a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
+++ b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
@@ -17,6 +17,7 @@
#include "mlir/Dialect/LLVMIR/FunctionCallUtils.h"
#include "mlir/Dialect/LLVMIR/LLVMAttrs.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/DialectResourceBlobManager.h"
#include "mlir/IR/TypeUtilities.h"
#include <type_traits>
@@ -430,10 +431,16 @@ static TypedAttr convertConstantValue(TypedAttr attr, Type resultType) {
retypeValues(cast<DenseIntElementsAttr>(sparseAttr.getValues())));
// A resource-backed elements attribute refers to a blob laid out for its own
- // element type, so it cannot be retyped here. Keep it rather than fail the
- // lowering.
- if (isa<ElementsAttr>(attr))
- return attr;
+ // element type. The blob cannot be rewritten here, only reinterpreted, which
+ // is correct exactly when the target type has the same width as the storage
+ // `index` uses in a blob.
+ if (auto resourceAttr = dyn_cast<DenseResourceElementsAttr>(attr)) {
+ if (width != IndexType::kInternalStorageBitWidth)
+ return {};
+ return DenseResourceElementsAttr::get(
+ cast<ShapedType>(attr.getType()).clone(targetIntType),
+ resourceAttr.getRawHandle());
+ }
return {};
}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
index a5c7a1b8c03b0..183cb6d69619e 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -3651,13 +3651,52 @@ LogicalResult LLVM::ConstantOp::verify() {
return success();
};
+ // Check that an integer attribute whose element type is `attributeIntType`
+ // is compatible with a type whose element type is `constantElementType`.
+ //
+ // Contrary to floats, integers must match exactly. An integer attribute
+ // carries no information that the corresponding LLVM type cannot represent,
+ // so any difference in width, signedness, or in `index` versus a fixed-width
+ // integer indicates a malformed constant. Note that `index` never reaches
+ // this check as a constant type, since it is not LLVM dialect-compatible.
+ auto verifyIntegerSemantics = [this](Type attributeIntType,
+ Type constantElementType,
+ StringRef description) -> LogicalResult {
+ if (attributeIntType != constantElementType)
+ return emitOpError() << "attribute and type have different integer "
+ << description << "s: " << attributeIntType
+ << " vs. " << constantElementType;
+ return success();
+ };
+
// Verification of IntegerAttr, FloatAttr, ElementsAttr, ArrayAttr.
- if (isa<IntegerAttr>(getValue())) {
+ if (auto intAttr = dyn_cast<IntegerAttr>(getValue())) {
if (!llvm::isa<IntegerType>(getType()))
return emitOpError() << "expected integer type";
+ return verifyIntegerSemantics(intAttr.getType(), getType(), "type");
} else if (auto floatAttr = dyn_cast<FloatAttr>(getValue())) {
return verifyFloatSemantics(floatAttr.getValue().getSemantics(), getType());
} else if (auto elementsAttr = dyn_cast<ElementsAttr>(getValue())) {
+ // Check that the element type of the attribute is compatible with the
+ // element type of the constant. Shared by the scalable and the fixed-size
+ // paths, since element types must agree either way.
+ auto verifyElementTypes = [&](ElementsAttr attr) -> LogicalResult {
+ Type attrElmType = LLVM::getConstantElementType(attr.getType());
+ Type resultElmType = LLVM::getConstantElementType(getType());
+ if (auto floatType = dyn_cast<FloatType>(attrElmType))
+ return verifyFloatSemantics(floatType.getFloatSemantics(),
+ resultElmType);
+
+ if (isa<IntegerType, IndexType>(attrElmType)) {
+ if (!isa<IntegerType>(resultElmType))
+ return emitOpError(
+ "expected integer element type for integer elements attribute");
+ return verifyIntegerSemantics(attrElmType, resultElmType,
+ "element type");
+ }
+ return success();
+ };
+
if (hasScalableVectorType(getType())) {
// The exact number of elements of a scalable vector is unknown, so we
// allow only splat attributes.
@@ -3665,7 +3704,7 @@ LogicalResult LLVM::ConstantOp::verify() {
if (!splatElementsAttr)
return emitOpError()
<< "scalable vector type requires a splat attribute";
- return success();
+ return verifyElementTypes(splatElementsAttr);
}
if (!isa<VectorType, LLVM::LLVMArrayType>(getType()))
return emitOpError() << "expected vector or array type";
@@ -3678,15 +3717,7 @@ LogicalResult LLVM::ConstantOp::verify() {
<< getNumElements(getType()) << " vs. " << attrNumElements;
}
- Type attrElmType = LLVM::getConstantElementType(elementsAttr.getType());
- Type resultElmType = LLVM::getConstantElementType(getType());
- if (auto floatType = dyn_cast<FloatType>(attrElmType))
- return verifyFloatSemantics(floatType.getFloatSemantics(), resultElmType);
-
- if (isa<IntegerType>(attrElmType) && !isa<IntegerType>(resultElmType)) {
- return emitOpError(
- "expected integer element type for integer elements attribute");
- }
+ return verifyElementTypes(elementsAttr);
} else if (auto arrayAttr = dyn_cast<ArrayAttr>(getValue())) {
// The case where the constant is LLVMStructType has already been handled.
diff --git a/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir b/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir
index 1341a6ca40019..c21e967976e4c 100644
--- a/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir
+++ b/mlir/test/Conversion/ArithToLLVM/arith-to-llvm.mlir
@@ -1068,10 +1068,11 @@ func.func @sparse_index_constant() -> vector<4xindex> {
// -----
// A resource-backed elements attribute refers to a blob laid out for its own
-// element type, so it is kept as is instead of being retyped.
+// element type, so it is reinterpreted rather than rewritten. This works because
+// `index` is stored with the same width as the target `i64`.
// CHECK-LABEL: @resource_index_constant
-// CHECK: llvm.mlir.constant(dense_resource<index_blob> : vector<2xindex>) : vector<2xi64>
+// CHECK: llvm.mlir.constant(dense_resource<index_blob> : vector<2xi64>) : vector<2xi64>
func.func @resource_index_constant() -> vector<2xindex> {
%0 = arith.constant dense_resource<index_blob> : vector<2xindex>
return %0 : vector<2xindex>
diff --git a/mlir/test/Conversion/ArithToLLVM/constant-index-bitwidth.mlir b/mlir/test/Conversion/ArithToLLVM/constant-index-bitwidth.mlir
index ddb7509b89369..4363286162bcc 100644
--- a/mlir/test/Conversion/ArithToLLVM/constant-index-bitwidth.mlir
+++ b/mlir/test/Conversion/ArithToLLVM/constant-index-bitwidth.mlir
@@ -42,3 +42,26 @@ func.func @sparse_index_constant() -> vector<4xindex> {
%0 = arith.constant sparse<[[0]], [-1]> : vector<4xindex>
return %0 : vector<4xindex>
}
+
+// A resource-backed elements attribute is only reinterpreted, so it can be
+// retyped exactly when the converted index type has the same width as the
+// storage `index` uses in a blob. Otherwise the lowering fails.
+
+// CHECK32-LABEL: @resource_index_constant
+// CHECK32: arith.constant dense_resource<index_blob> : vector<2xindex>
+// CHECK64-LABEL: @resource_index_constant
+// CHECK64: llvm.mlir.constant(dense_resource<index_blob> : vector<2xi64>) : vector<2xi64>
+// CHECK128-LABEL: @resource_index_constant
+// CHECK128: arith.constant dense_resource<index_blob> : vector<2xindex>
+func.func @resource_index_constant() -> vector<2xindex> {
+ %0 = arith.constant dense_resource<index_blob> : vector<2xindex>
+ return %0 : vector<2xindex>
+}
+
+{-#
+ dialect_resources: {
+ builtin: {
+ index_blob: "0x0800000001000000000000000200000000000000"
+ }
+ }
+#-}
diff --git a/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir b/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir
index 1594e18077196..6445f5afd8000 100644
--- a/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir
+++ b/mlir/test/Conversion/MPIToLLVM/mpitollvm.mlir
@@ -341,7 +341,7 @@ module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">,
// CHECK: [[v1:%.*]] = llvm.extractvalue [[v0]][1] : !llvm.struct<(ptr, ptr, i32, array<1 x i32>, array<1 x i32>)>
// CHECK: [[v2:%.*]] = llvm.extractvalue [[v0]][2] : !llvm.struct<(ptr, ptr, i32, array<1 x i32>, array<1 x i32>)>
// CHECK: [[v3:%.*]] = llvm.getelementptr [[v1]][[[v2]]] : (!llvm.ptr, i32) -> !llvm.ptr, f32
- // CHECK: [[v4:%.*]] = llvm.mlir.constant(1 : index) : i32
+ // CHECK: [[v4:%.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: [[v5:%.*]] = llvm.extractvalue [[v0]][3, 0] : !llvm.struct<(ptr, ptr, i32, array<1 x i32>, array<1 x i32>)>
// COM: No width adjustment, the extent is already an `i32`.
// CHECK-NOT: llvm.trunc
@@ -366,7 +366,7 @@ module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">,
// CHECK: [[v1:%.*]] = llvm.extractvalue [[v0]][1] : !llvm.struct<(ptr, ptr, i16, array<1 x i16>, array<1 x i16>)>
// CHECK: [[v2:%.*]] = llvm.extractvalue [[v0]][2] : !llvm.struct<(ptr, ptr, i16, array<1 x i16>, array<1 x i16>)>
// CHECK: [[v3:%.*]] = llvm.getelementptr [[v1]][[[v2]]] : (!llvm.ptr, i16) -> !llvm.ptr, f32
- // CHECK: [[v4:%.*]] = llvm.mlir.constant(1 : index) : i32
+ // CHECK: [[v4:%.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: [[v5:%.*]] = llvm.extractvalue [[v0]][3, 0] : !llvm.struct<(ptr, ptr, i16, array<1 x i16>, array<1 x i16>)>
// CHECK: [[v6:%.*]] = llvm.zext [[v5]] : i16 to i32
// CHECK: [[v7:%.*]] = llvm.mul [[v6]], [[v4]] : i32
@@ -388,7 +388,7 @@ module attributes {dlti.map = #dlti.map<"MPI:Implementation" = "MPICH">} {
// CHECK: [[v1:%.*]] = llvm.extractvalue [[v0]][1] : !llvm.struct<(ptr, ptr, i64)>
// CHECK: [[v2:%.*]] = llvm.extractvalue [[v0]][2] : !llvm.struct<(ptr, ptr, i64)>
// CHECK: [[v3:%.*]] = llvm.getelementptr [[v1]][[[v2]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
- // CHECK: [[v4:%.*]] = llvm.mlir.constant(1 : index) : i32
+ // CHECK: [[v4:%.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK-NOT: llvm.mul
// CHECK: llvm.call @MPI_Send([[v3]], [[v4]], {{.*}}) : (!llvm.ptr, i32, i32, i32, i32, i32) -> i32
mpi.send(%arg0, %rank, %rank, %comm) : memref<f32>, i32, i32
diff --git a/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir b/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir
index 65397f7cbdff4..13607ee111518 100644
--- a/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir
+++ b/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir
@@ -330,7 +330,7 @@ llvm.func @_QPsb() {
// CHECK: omp.yield(%[[TRUE_EXT]] : i32)
// CHECK: } combiner {
// CHECK: ^bb0(%[[ARG_1:.*]]: i32, %[[ARG_2:.*]]: i32):
-// CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i64) : i32
+// CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
// CHECK: %[[CMP_1:.*]] = llvm.icmp "ne" %[[ARG_1]], %[[ZERO]] : i32
// CHECK: %[[CMP_2:.*]] = llvm.icmp "ne" %[[ARG_2]], %[[ZERO]] : i32
// CHECK: %[[COMBINE_VAL:.*]] = llvm.icmp "eq" %[[CMP_1]], %[[CMP_2]] : i1
@@ -356,7 +356,7 @@ omp.declare_reduction @eqv_reduction : i32 init {
omp.yield(%1 : i32)
} combiner {
^bb0(%arg0: i32, %arg1: i32):
- %0 = llvm.mlir.constant(0 : i64) : i32
+ %0 = llvm.mlir.constant(0 : i32) : i32
%1 = llvm.icmp "ne" %arg0, %0 : i32
%2 = llvm.icmp "ne" %arg1, %0 : i32
%3 = llvm.icmp "eq" %1, %2 : i1
diff --git a/mlir/test/Dialect/LLVMIR/inlining-alias-scopes.mlir b/mlir/test/Dialect/LLVMIR/inlining-alias-scopes.mlir
index 6b369c5012105..aea20df9c6de0 100644
--- a/mlir/test/Dialect/LLVMIR/inlining-alias-scopes.mlir
+++ b/mlir/test/Dialect/LLVMIR/inlining-alias-scopes.mlir
@@ -264,7 +264,7 @@ llvm.func @missing_noalias_on_one_ptr(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2:
llvm.func @might_return_arg_derived(!llvm.ptr) -> !llvm.ptr
llvm.func @foo(%arg0: !llvm.ptr {llvm.noalias}, %arg1: !llvm.ptr {llvm.noalias}) {
- %0 = llvm.mlir.constant(5 : i64) : i32
+ %0 = llvm.mlir.constant(5 : i32) : i32
%1 = llvm.call @might_return_arg_derived(%arg0) : (!llvm.ptr) -> !llvm.ptr
llvm.store %0, %1 : i32, !llvm.ptr
llvm.return
@@ -291,7 +291,7 @@ llvm.func @bar(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr) {
llvm.func @random() -> i1
llvm.func @block_arg(%arg0: !llvm.ptr {llvm.noalias}, %arg1: !llvm.ptr {llvm.noalias}) {
- %0 = llvm.mlir.constant(5 : i64) : i32
+ %0 = llvm.mlir.constant(5 : i32) : i32
%1 = llvm.call @random() : () -> i1
llvm.cond_br %1, ^bb0(%arg0 : !llvm.ptr), ^bb0(%arg1 : !llvm.ptr)
@@ -320,7 +320,7 @@ llvm.func @bar(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr) {
llvm.func @random() -> i1
llvm.func @region_branch(%arg0: !llvm.ptr {llvm.noalias}, %arg1: !llvm.ptr {llvm.noalias}) {
- %0 = llvm.mlir.constant(5 : i64) : i32
+ %0 = llvm.mlir.constant(5 : i32) : i32
test.region_if %arg0: !llvm.ptr -> !llvm.ptr then {
^bb0(%arg2: !llvm.ptr):
test.region_if_yield %arg0 : !llvm.ptr
@@ -347,7 +347,7 @@ llvm.func @region_branch_inlining(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !ll
// -----
llvm.func @missing_region_branch(%arg0: !llvm.ptr {llvm.noalias}, %arg1: !llvm.ptr {llvm.noalias}) {
- %0 = llvm.mlir.constant(5 : i64) : i32
+ %0 = llvm.mlir.constant(5 : i32) : i32
"test.one_region_op"() ({
^bb0(%arg2: !llvm.ptr):
llvm.store %0, %arg2 : i32, !llvm.ptr
@@ -374,7 +374,7 @@ llvm.func @missing_region_branch_inlining(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %a
llvm.func @random() -> i1
llvm.func @block_arg(%arg0: !llvm.ptr {llvm.noalias}, %arg1: !llvm.ptr {llvm.noalias}) {
- %0 = llvm.mlir.constant(5 : i64) : i32
+ %0 = llvm.mlir.constant(5 : i32) : i32
%1 = llvm.mlir.constant(1 : i64) : i64
%2 = llvm.alloca %1 x i32 : (i64) -> !llvm.ptr
%3 = llvm.call @random() : () -> i1
@@ -407,7 +407,7 @@ llvm.func @unknown() -> !llvm.ptr
llvm.func @random() -> i1
llvm.func @unknown_object(%arg0: !llvm.ptr {llvm.noalias}, %arg1: !llvm.ptr {llvm.noalias}) {
- %0 = llvm.mlir.constant(5 : i64) : i32
+ %0 = llvm.mlir.constant(5 : i32) : i32
%1 = llvm.call @random() : () -> i1
%2 = llvm.call @unknown() : () -> !llvm.ptr
llvm.cond_br %1, ^bb0(%arg0 : !llvm.ptr), ^bb0(%2 : !llvm.ptr)
@@ -439,7 +439,7 @@ llvm.func @bar(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr) {
// CHECK-DAG: #[[$ARG1_SCOPE:.*]] = #llvm.alias_scope<id = {{.*}}, domain = #[[DOMAIN]]{{(,.*)?}}>
llvm.func @supported_operations(%arg0: !llvm.ptr {llvm.noalias}, %arg1: !llvm.ptr {llvm.noalias}) {
- %0 = llvm.mlir.constant(5 : i64) : i32
+ %0 = llvm.mlir.constant(5 : i32) : i32
llvm.store %0, %arg1 : i32, !llvm.ptr
%1 = llvm.load %arg1 : !llvm.ptr -> i32
"llvm.intr.memcpy"(%arg0, %arg1, %1) <{ isVolatile = false }> : (!llvm.ptr, !llvm.ptr, i32) -> ()
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index b61856b422876..7ef72ddf709c0 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -472,6 +472,62 @@ llvm.func @vector_int_attr_requires_int_type() -> vector<2xf32> {
// -----
+llvm.func @int_attr_and_type_required_same_index() -> i64 {
+ // expected-error @below{{attribute and type have different integer types: 'index' vs. 'i64'}}
+ %0 = llvm.mlir.constant(1 : index) : i64
+ llvm.return %0 : i64
+}
+
+// -----
+
+llvm.func @int_attr_and_type_required_same_width() -> i16 {
+ // expected-error @below{{attribute and type have different integer types: 'i8' vs. 'i16'}}
+ %0 = llvm.mlir.constant(1 : i8) : i16
+ llvm.return %0 : i16
+}
+
+// -----
+
+llvm.func @int_attr_and_type_required_same_signedness() -> i32 {
+ // expected-error @below{{attribute and type have different integer types: 'ui32' vs. 'i32'}}
+ %0 = llvm.mlir.constant(1 : ui32) : i32
+ llvm.return %0 : i32
+}
+
+// -----
+
+llvm.func @vector_int_attr_and_type_required_same_index() -> vector<2xi64> {
+ // expected-error @below{{attribute and type have different integer element types: 'index' vs. 'i64'}}
+ %0 = llvm.mlir.constant(dense<[1, 2]> : vector<2xindex>) : vector<2xi64>
+ llvm.return %0 : vector<2xi64>
+}
+
+// -----
+
+llvm.func @vector_int_attr_and_type_required_same_width() -> vector<2xi64> {
+ // expected-error @below{{attribute and type have different integer element types: 'i32' vs. 'i64'}}
+ %0 = llvm.mlir.constant(dense<[1, 2]> : vector<2xi32>) : vector<2xi64>
+ llvm.return %0 : vector<2xi64>
+}
+
+// -----
+
+llvm.func @scalable_vector_int_attr_and_type_required_same() -> vector<[2]xi64> {
+ // expected-error @below{{attribute and type have different integer element types: 'i32' vs. 'i64'}}
+ %0 = llvm.mlir.constant(dense<1> : vector<2xi32>) : vector<[2]xi64>
+ llvm.return %0 : vector<[2]xi64>
+}
+
+// -----
+
+llvm.func @array_int_attr_and_type_required_same() -> !llvm.array<2 x i64> {
+ // expected-error @below{{attribute and type have different integer element types: 'i32' vs. 'i64'}}
+ %0 = llvm.mlir.constant(dense<[1, 2]> : tensor<2xi32>) : !llvm.array<2 x i64>
+ llvm.return %0 : !llvm.array<2 x i64>
+}
+
+// -----
+
llvm.func @float_attr_and_type_required_same() -> f16 {
// expected-error @below{{attribute and type have different float semantics}}
%cst = llvm.mlir.constant(1.0 : bf16) : f16
diff --git a/mlir/test/Dialect/LLVMIR/roundtrip.mlir b/mlir/test/Dialect/LLVMIR/roundtrip.mlir
index e8490d55f2c05..b27d07ecdb4f5 100644
--- a/mlir/test/Dialect/LLVMIR/roundtrip.mlir
+++ b/mlir/test/Dialect/LLVMIR/roundtrip.mlir
@@ -198,9 +198,9 @@ func.func @ops(%arg0: i32, %arg1: f32,
// CHECK: ^[[BB2]]
^bb2:
// CHECK: %{{.*}} = llvm.mlir.undef : !llvm.struct<(i32, f64, i32)>
-// CHECK: %{{.*}} = llvm.mlir.constant(42 : i64) : i47
+// CHECK: %{{.*}} = llvm.mlir.constant(42 : i47) : i47
%22 = llvm.mlir.undef : !llvm.struct<(i32, f64, i32)>
- %23 = llvm.mlir.constant(42) : i47
+ %23 = llvm.mlir.constant(42 : i47) : i47
// CHECK: llvm.switch %0 : i32, ^[[BB3]] [
// CHECK-NEXT: 1: ^[[BB4:.*]],
// CHECK-NEXT: 2: ^[[BB5:.*]],
@@ -318,25 +318,25 @@ llvm.func @vararg_foo(i32, ...) -> !llvm.struct<(i32, f64, i32)>
// An larger self-contained function.
// CHECK-LABEL: llvm.func @foo(%{{.*}}: i32) -> !llvm.struct<(i32, f64, i32)> {
llvm.func @foo(%arg0: i32) -> !llvm.struct<(i32, f64, i32)> {
-// CHECK: %[[V0:.*]] = llvm.mlir.constant(3 : i64) : i32
-// CHECK: %[[V1:.*]] = llvm.mlir.constant(3 : i64) : i32
+// CHECK: %[[V0:.*]] = llvm.mlir.constant(3 : i32) : i32
+// CHECK: %[[V1:.*]] = llvm.mlir.constant(3 : i32) : i32
// CHECK: %[[V2:.*]] = llvm.mlir.constant(4.200000e+01 : f64) : f64
// CHECK: %[[V3:.*]] = llvm.mlir.constant(4.200000e+01 : f64) : f64
// CHECK: %[[V4:.*]] = llvm.add %[[V0]], %[[V1]] : i32
// CHECK: %[[V5:.*]] = llvm.mul %[[V4]], %[[V1]] : i32
// CHECK: %[[V6:.*]] = llvm.fadd %[[V2]], %[[V3]] : f64
// CHECK: %[[V7:.*]] = llvm.fsub %[[V3]], %[[V6]] : f64
-// CHECK: %[[V8:.*]] = llvm.mlir.constant(1 : i64) : i1
+// CHECK: %[[V8:.*]] = llvm.mlir.constant(true) : i1
// CHECK: llvm.cond_br %[[V8]], ^[[BB1:.*]](%[[V4]] : i32), ^[[BB2:.*]](%[[V4]] : i32)
- %0 = llvm.mlir.constant(3) : i32
- %1 = llvm.mlir.constant(3) : i32
+ %0 = llvm.mlir.constant(3 : i32) : i32
+ %1 = llvm.mlir.constant(3 : i32) : i32
%2 = llvm.mlir.constant(4.200000e+01) : f64
%3 = llvm.mlir.constant(4.200000e+01) : f64
%4 = llvm.add %0, %1 : i32
%5 = llvm.mul %4, %1 : i32
%6 = llvm.fadd %2, %3 : f64
%7 = llvm.fsub %3, %6 : f64
- %8 = llvm.mlir.constant(1) : i1
+ %8 = llvm.mlir.constant(true) : i1
llvm.cond_br %8, ^bb1(%4 : i32), ^bb2(%4 : i32)
// CHECK:^[[BB1]](%[[V9:.*]]: i32):
@@ -1049,7 +1049,7 @@ llvm.func @test_invoke_with_opbundle() attributes { personality = @__gxx_persona
// CHECK-LABEL: @test_call_intrin_with_opbundle
llvm.func @test_call_intrin_with_opbundle(%arg0 : !llvm.ptr) {
- %0 = llvm.mlir.constant(1 : i1) : i1
+ %0 = llvm.mlir.constant(true) : i1
%1 = llvm.mlir.constant(16 : i32) : i32
// CHECK: llvm.call_intrinsic "llvm.assume"(%{{.+}}) ["align"(%{{.+}}, %{{.+}} : !llvm.ptr, i32)] : (i1) -> ()
llvm.call_intrinsic "llvm.assume"(%0) ["align"(%arg0, %1 : !llvm.ptr, i32)] : (i1) -> ()
@@ -1058,7 +1058,7 @@ llvm.func @test_call_intrin_with_opbundle(%arg0 : !llvm.ptr) {
// CHECK-LABEL: @test_assume_intr_no_opbundle
llvm.func @test_assume_intr_no_opbundle(%arg0 : !llvm.ptr) {
- %0 = llvm.mlir.constant(1 : i1) : i1
+ %0 = llvm.mlir.constant(true) : i1
// CHECK: llvm.intr.assume %0 : i1
llvm.intr.assume %0 : i1
llvm.return
@@ -1066,7 +1066,7 @@ llvm.func @test_assume_intr_no_opbundle(%arg0 : !llvm.ptr) {
// CHECK-LABEL: @test_assume_intr_empty_opbundle
llvm.func @test_assume_intr_empty_opbundle(%arg0 : !llvm.ptr) {
- %0 = llvm.mlir.constant(1 : i1) : i1
+ %0 = llvm.mlir.constant(true) : i1
// CHECK: llvm.intr.assume %0 : i1
llvm.intr.assume %0 [] : i1
llvm.return
@@ -1074,7 +1074,7 @@ llvm.func @test_assume_intr_empty_opbundle(%arg0 : !llvm.ptr) {
// CHECK-LABEL: @test_assume_intr_with_opbundles
llvm.func @test_assume_intr_with_opbundles(%arg0 : !llvm.ptr) {
- %0 = llvm.mlir.constant(1 : i1) : i1
+ %0 = llvm.mlir.constant(true) : i1
%1 = llvm.mlir.constant(2 : i32) : i32
%2 = llvm.mlir.constant(3 : i32) : i32
%3 = llvm.mlir.constant(4 : i32) : i32
diff --git a/mlir/test/Dialect/Vector/canonicalize.mlir b/mlir/test/Dialect/Vector/canonicalize.mlir
index 7ad6eda6ec1ba..f6ba206e7449e 100644
--- a/mlir/test/Dialect/Vector/canonicalize.mlir
+++ b/mlir/test/Dialect/Vector/canonicalize.mlir
@@ -3842,7 +3842,7 @@ func.func @from_elements_f64_to_i64_conversion() -> vector<6xi64> {
// CHECK-NEXT: %[[CST:.*]] = arith.constant dense<0> : vector<1xi8>
// CHECK-NEXT: return %[[CST]] : vector<1xi8>
func.func @from_elements_i1_to_i8_conversion() -> vector<1xi8> {
- %cst = llvm.mlir.constant(0: i1) : i8
+ %cst = llvm.mlir.constant(0: i8) : i8
%v = vector.from_elements %cst : vector<1xi8>
return %v : vector<1xi8>
}
diff --git a/mlir/test/Target/LLVMIR/nvvm/fence.mlir b/mlir/test/Target/LLVMIR/nvvm/fence.mlir
index b2e81d9586870..423a6e97d1d4a 100644
--- a/mlir/test/Target/LLVMIR/nvvm/fence.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/fence.mlir
@@ -74,7 +74,7 @@ llvm.func @nvvm_fence_proxy_tensormap_generic_release() {
// CHECK-LABEL: @nvvm_fence_proxy_tensormap_generic_acquire
llvm.func @nvvm_fence_proxy_tensormap_generic_acquire(%addr : !llvm.ptr) {
- %c128 = llvm.mlir.constant(128) : i32
+ %c128 = llvm.mlir.constant(128 : i32) : i32
// CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr {{%[0-9]+}}, i32 128)
nvvm.fence.proxy.acquire cta %addr, %c128
diff --git a/mlir/test/Target/LLVMIR/omptarget-if-nowait.mlir b/mlir/test/Target/LLVMIR/omptarget-if-nowait.mlir
index e055c2f7324c7..2b5e349aff2e2 100644
--- a/mlir/test/Target/LLVMIR/omptarget-if-nowait.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-if-nowait.mlir
@@ -5,7 +5,7 @@ module attributes {omp.is_target_device = false, omp.target_triples = ["amdgcn-a
%0 = llvm.mlir.constant(1 : i64) : i64
%3 = llvm.alloca %0 x i32 {bindc_name = "cond"} : (i64) -> !llvm.ptr
%6 = llvm.load %3 : !llvm.ptr -> i32
- %7 = llvm.mlir.constant(0 : i64) : i32
+ %7 = llvm.mlir.constant(0 : i32) : i32
%8 = llvm.icmp "ne" %6, %7 : i32
%9 = omp.map.info var_ptr(%3 : !llvm.ptr, i32) map_clauses(implicit, exit_release_or_enter_alloc) capture(ByCopy) name("cond") -> !llvm.ptr
%10 = omp.map.info var_ptr(%arg0 : !llvm.ptr, f32) map_clauses(implicit, exit_release_or_enter_alloc) capture(ByCopy) name("var") -> !llvm.ptr
diff --git a/mlir/test/Target/LLVMIR/omptarget-parallel-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-parallel-llvm.mlir
index b45016b612656..d05b2c4c06173 100644
--- a/mlir/test/Target/LLVMIR/omptarget-parallel-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-parallel-llvm.mlir
@@ -40,7 +40,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memo
omp.target kernel_type(generic) map_entries(%2 -> %arg1, %3 -> %arg2 : !llvm.ptr, !llvm.ptr) {
%4 = llvm.mlir.constant(10 : i32) : i32
%5 = llvm.load %arg2 : !llvm.ptr -> i32
- %6 = llvm.mlir.constant(0 : i64) : i32
+ %6 = llvm.mlir.constant(0 : i32) : i32
%7 = llvm.icmp "ne" %5, %6 : i32
omp.parallel if(%7) {
llvm.store %4, %arg1 : i32, !llvm.ptr
diff --git a/mlir/test/Target/LLVMIR/openmp-composite-simd-if.mlir b/mlir/test/Target/LLVMIR/openmp-composite-simd-if.mlir
index f8ebe6c9ba16f..2f82bb68af824 100644
--- a/mlir/test/Target/LLVMIR/openmp-composite-simd-if.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-composite-simd-if.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
llvm.func @_QPfoo(%arg0: !llvm.ptr {fir.bindc_name = "array", llvm.nocapture}, %arg1: !llvm.ptr {fir.bindc_name = "t", llvm.nocapture}) {
- %0 = llvm.mlir.constant(0 : i64) : i32
+ %0 = llvm.mlir.constant(0 : i32) : i32
%1 = llvm.mlir.constant(1 : i32) : i32
%2 = llvm.mlir.constant(10 : i64) : i64
%3 = llvm.mlir.constant(1 : i64) : i64
diff --git a/mlir/test/Target/LLVMIR/openmp-llvm.mlir b/mlir/test/Target/LLVMIR/openmp-llvm.mlir
index ce237da85906d..7be951a8f49c9 100644
--- a/mlir/test/Target/LLVMIR/openmp-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-llvm.mlir
@@ -3686,7 +3686,7 @@ llvm.func @bar(i32, i32, !llvm.ptr) -> ()
llvm.func @omp_taskgroup_task(%x: i32, %y: i32, %zaddr: !llvm.ptr) {
omp.taskgroup {
- %c1 = llvm.mlir.constant(1) : i32
+ %c1 = llvm.mlir.constant(1 : i32) : i32
%ptr1 = llvm.alloca %c1 x i8 : (i32) -> !llvm.ptr
omp.task {
llvm.call @foo() : () -> ()
diff --git a/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir b/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
index e8ca7567d3ae2..624ea90d4c8ce 100644
--- a/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
@@ -16,10 +16,10 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memo
llvm.func @main(%num_teams : !llvm.ptr) {
// CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_main_l{{[0-9]+}}(ptr %[[NUM_TEAMS_ARG:.*]], ptr %[[KERNEL_ARGS:.*]]) #[[ATTRS1:[0-9]+]]
// CHECK: %{{.*}} = call i32 @__kmpc_target_init(ptr @[[KERNEL1_ENV]], ptr %[[KERNEL_ARGS]])
- %target_threads = llvm.mlir.constant(20) : i32
+ %target_threads = llvm.mlir.constant(20 : i32) : i32
%0 = omp.map.info var_ptr(%num_teams : !llvm.ptr, i32) map_clauses(to) capture(ByCopy) -> !llvm.ptr
omp.target kernel_type(generic) thread_limit(%target_threads : i32) map_entries(%0 -> %arg_teams : !llvm.ptr) {
- %teams_threads = llvm.mlir.constant(10) : i32
+ %teams_threads = llvm.mlir.constant(10 : i32) : i32
%num_teams1 = llvm.load %arg_teams : !llvm.ptr -> i32
omp.teams num_teams(to %num_teams1 : i32) thread_limit(%teams_threads : i32) {
omp.terminator
@@ -29,9 +29,9 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memo
// CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_main_l{{[0-9]+}}(ptr %[[KERNEL_ARGS:.*]]) #[[ATTRS2:[0-9]+]]
// CHECK: %{{.*}} = call i32 @__kmpc_target_init(ptr @[[KERNEL2_ENV]], ptr %[[KERNEL_ARGS]])
- %target_threads2 = llvm.mlir.constant(30) : i32
+ %target_threads2 = llvm.mlir.constant(30 : i32) : i32
omp.target kernel_type(generic) thread_limit(%target_threads2 : i32) {
- %num_teams2 = llvm.mlir.constant(40) : i32
+ %num_teams2 = llvm.mlir.constant(40 : i32) : i32
omp.teams num_teams(to %num_teams2 : i32) {
omp.terminator
}
diff --git a/mlir/test/Target/LLVMIR/openmp-target-launch-host.mlir b/mlir/test/Target/LLVMIR/openmp-target-launch-host.mlir
index 67f8c7c3eac0f..9521dbf9d8e5a 100644
--- a/mlir/test/Target/LLVMIR/openmp-target-launch-host.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-target-launch-host.mlir
@@ -17,8 +17,8 @@
// CHECK: call void @__kmpc_push_num_teams_51(ptr {{.*}}, i32 {{.*}}, i32 %[[NUM_TEAMS_OUTLINED]], i32 %[[NUM_TEAMS_OUTLINED]], i32 [[NUM_THREADS]])
module attributes {omp.is_target_device = false, omp.target_triples = ["amdgcn-amd-amdhsa"]} {
llvm.func @main(%num_teams : i32) {
- %target_threads = llvm.mlir.constant(20) : i32
- %teams_threads = llvm.mlir.constant(10) : i32
+ %target_threads = llvm.mlir.constant(20 : i32) : i32
+ %teams_threads = llvm.mlir.constant(10 : i32) : i32
omp.target kernel_type(generic) thread_limit(%target_threads : i32)
host_eval(%num_teams -> %arg_teams, %teams_threads -> %arg_teams_threads : i32, i32) {
omp.teams num_teams(to %arg_teams : i32) thread_limit(%arg_teams_threads : i32) {
diff --git a/mlir/test/Transforms/sccp.mlir b/mlir/test/Transforms/sccp.mlir
index 80273c86090a7..a12e7dc9cc99c 100644
--- a/mlir/test/Transforms/sccp.mlir
+++ b/mlir/test/Transforms/sccp.mlir
@@ -247,12 +247,15 @@ func.func @op_with_region() -> (i32) {
return %1 : i32
}
+// `llvm.mlir.constant` allows a float attribute to be paired with an integer
+// type of the same width, so the folded attribute type differs from the source
+// type of the broadcast. Folding must bail out instead of crashing.
// CHECK-LABEL: no_crash_with_different_source_type
func.func @no_crash_with_different_source_type() {
- // CHECK: llvm.mlir.constant(0 : index) : i64
- %0 = llvm.mlir.constant(0 : index) : i64
- // CHECK: vector.broadcast %[[CST:.*]] : i64 to vector<128xi64>
- %1 = vector.broadcast %0 : i64 to vector<128xi64>
+ // CHECK: llvm.mlir.constant(0.000000e+00 : f32) : i32
+ %0 = llvm.mlir.constant(0.0 : f32) : i32
+ // CHECK: vector.broadcast %[[CST:.*]] : i32 to vector<128xi32>
+ %1 = vector.broadcast %0 : i32 to vector<128xi32>
llvm.return
}
More information about the llvm-branch-commits
mailing list