[Mlir-commits] [mlir] [mlir][vector] Add alignment attribute to vector operations. (PR #152507)
Erick Ochoa Lopez
llvmlistbot at llvm.org
Mon Aug 11 06:39:50 PDT 2025
https://github.com/amd-eochoalo updated https://github.com/llvm/llvm-project/pull/152507
>From a97cbc6a1223fd34dc9e12e6fe8f4d3445fac47e Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Wed, 6 Aug 2025 19:48:16 -0700
Subject: [PATCH 01/10] [mlir][vector] Add alignment to vector.gather.
---
.../mlir/Dialect/Vector/IR/VectorOps.td | 29 ++++++++++++++++++-
mlir/test/Dialect/Vector/invalid.mlir | 18 ++++++++++++
2 files changed, 46 insertions(+), 1 deletion(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index b3b8afdd8b4c1..250bd55e62377 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -2054,7 +2054,9 @@ def Vector_GatherOp :
Variadic<Index>:$indices,
VectorOfNonZeroRankOf<[AnyInteger, Index]>:$index_vec,
VectorOfNonZeroRankOf<[I1]>:$mask,
- AnyVectorOfNonZeroRank:$pass_thru)>,
+ AnyVectorOfNonZeroRank:$pass_thru,
+ ConfinedAttr<OptionalAttr<I64Attr>,
+ [AllAttrOf<[IntPositive, IntPowerOf2]>]>:$alignment)>,
Results<(outs AnyVectorOfNonZeroRank:$result)> {
let summary = [{
@@ -2111,6 +2113,31 @@ def Vector_GatherOp :
"`into` type($result)";
let hasCanonicalizer = 1;
let hasVerifier = 1;
+
+ let builders = [
+ OpBuilder<(ins "VectorType":$resultType,
+ "Value":$base,
+ "ValueRange":$indices,
+ "Value":$index_vec,
+ "Value":$mask,
+ "Value":$passthrough,
+ CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ return build($_builder, $_state, resultType, base, indices, index_vec, mask, passthrough,
+ alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ nullptr);
+ }]>,
+ OpBuilder<(ins "TypeRange":$resultTypes,
+ "Value":$base,
+ "ValueRange":$indices,
+ "Value":$index_vec,
+ "Value":$mask,
+ "Value":$passthrough,
+ CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ return build($_builder, $_state, resultTypes, base, indices, index_vec, mask, passthrough,
+ alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ nullptr);
+ }]>
+ ];
}
def Vector_ScatterOp :
diff --git a/mlir/test/Dialect/Vector/invalid.mlir b/mlir/test/Dialect/Vector/invalid.mlir
index 211e16db85a94..68b67a58bf736 100644
--- a/mlir/test/Dialect/Vector/invalid.mlir
+++ b/mlir/test/Dialect/Vector/invalid.mlir
@@ -1470,6 +1470,24 @@ func.func @gather_pass_thru_type_mismatch(%base: memref<?xf32>, %indices: vector
// -----
+func.func @gather_invalid_alignment(%base: memref<16xf32>, %indices: vector<16xi32>,
+ %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0 : index) {
+ // expected-error at +2 {{'vector.gather' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ %0 = vector.gather %base[%c0][%indices], %mask, %pass_thru
+ { alignment = -1 } : memref<16xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32> into vector<16xf32>
+}
+
+// -----
+
+func.func @gather_invalid_alignment(%base: memref<16xf32>, %indices: vector<16xi32>,
+ %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0 : index) {
+ // expected-error at +2 {{'vector.gather' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ %0 = vector.gather %base[%c0][%indices], %mask, %pass_thru
+ { alignment = 3 } : memref<16xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32> into vector<16xf32>
+}
+
+// -----
+
func.func @scatter_to_vector(%base: vector<16xf32>, %indices: vector<16xi32>,
%mask: vector<16xi1>, %pass_thru: vector<16xf32>) {
%c0 = arith.constant 0 : index
>From 158b91cf695cd3ad9b6893428666938daab6c86a Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Wed, 6 Aug 2025 19:58:43 -0700
Subject: [PATCH 02/10] [mlir][vector] Add alignment to vector.scatter
---
.../mlir/Dialect/Vector/IR/VectorOps.td | 17 ++++++++++++++++-
mlir/test/Dialect/Vector/invalid.mlir | 18 ++++++++++++++++++
2 files changed, 34 insertions(+), 1 deletion(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index 250bd55e62377..5dd452b8efd81 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -2146,7 +2146,9 @@ def Vector_ScatterOp :
Variadic<Index>:$indices,
VectorOfNonZeroRankOf<[AnyInteger, Index]>:$index_vec,
VectorOfNonZeroRankOf<[I1]>:$mask,
- AnyVectorOfNonZeroRank:$valueToStore)> {
+ AnyVectorOfNonZeroRank:$valueToStore,
+ ConfinedAttr<OptionalAttr<I64Attr>,
+ [AllAttrOf<[IntPositive, IntPowerOf2]>]>:$alignment)> {
let summary = [{
scatters elements from a vector into memory as defined by an index vector
@@ -2204,6 +2206,19 @@ def Vector_ScatterOp :
"type($index_vec) `,` type($mask) `,` type($valueToStore)";
let hasCanonicalizer = 1;
let hasVerifier = 1;
+
+ let builders = [
+ OpBuilder<(ins "Value":$base,
+ "ValueRange":$indices,
+ "Value":$index_vec,
+ "Value":$mask,
+ "Value":$valueToStore,
+ CArg<"llvm::Align", "llvm::Align()">: $alignment), [{
+ return build($_builder, $_state, base, indices, index_vec, mask, valueToStore,
+ alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ nullptr);
+ }]>
+ ];
}
def Vector_ExpandLoadOp :
diff --git a/mlir/test/Dialect/Vector/invalid.mlir b/mlir/test/Dialect/Vector/invalid.mlir
index 68b67a58bf736..dcc4c75c72595 100644
--- a/mlir/test/Dialect/Vector/invalid.mlir
+++ b/mlir/test/Dialect/Vector/invalid.mlir
@@ -1549,6 +1549,24 @@ func.func @scatter_dim_mask_mismatch(%base: memref<?xf32>, %indices: vector<16xi
// -----
+func.func @scatter_invalid_alignment(%base: memref<?xf32>, %indices: vector<16xi32>,
+ %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
+ // expected-error at +1 {{'vector.scatter' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ vector.scatter %base[%c0][%indices], %mask, %value { alignment = -1 }
+ : memref<?xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32>
+}
+
+// -----
+
+func.func @scatter_invalid_alignment(%base: memref<?xf32>, %indices: vector<16xi32>,
+ %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
+ // expected-error at +1 {{'vector.scatter' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ vector.scatter %base[%c0][%indices], %mask, %value { alignment = 3 }
+ : memref<?xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32>
+}
+
+// -----
+
func.func @expand_base_type_mismatch(%base: memref<?xf64>, %mask: vector<16xi1>, %pass_thru: vector<16xf32>) {
%c0 = arith.constant 0 : index
// expected-error at +1 {{'vector.expandload' op base and result element type should match}}
>From 36949d1cbce78fc518619c759c557afd1feccd3f Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Wed, 6 Aug 2025 20:48:21 -0700
Subject: [PATCH 03/10] [mlir][vector] Add alignment to compressstore
---
mlir/include/mlir/Dialect/Vector/IR/VectorOps.td | 15 ++++++++++++++-
mlir/test/Dialect/Vector/invalid.mlir | 14 ++++++++++++++
2 files changed, 28 insertions(+), 1 deletion(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index 5dd452b8efd81..0a36a54562fff 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -2295,7 +2295,9 @@ def Vector_CompressStoreOp :
Arguments<(ins Arg<AnyMemRef, "", [MemWrite]>:$base,
Variadic<Index>:$indices,
FixedVectorOfNonZeroRankOf<[I1]>:$mask,
- AnyVectorOfNonZeroRank:$valueToStore)> {
+ AnyVectorOfNonZeroRank:$valueToStore,
+ ConfinedAttr<OptionalAttr<I64Attr>,
+ [AllAttrOf<[IntPositive, IntPowerOf2]>]>:$alignment)> {
let summary = "writes elements selectively from a vector as defined by a mask";
@@ -2354,6 +2356,17 @@ def Vector_CompressStoreOp :
"type($base) `,` type($mask) `,` type($valueToStore)";
let hasCanonicalizer = 1;
let hasVerifier = 1;
+ let builders = [
+ OpBuilder<(ins "Value":$base,
+ "ValueRange":$indices,
+ "Value":$mask,
+ "Value":$valueToStore,
+ CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ return build($_builder, $_state, base, indices, valueToStore, mask,
+ alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ nullptr);
+ }]>
+ ];
}
def Vector_ShapeCastOp :
diff --git a/mlir/test/Dialect/Vector/invalid.mlir b/mlir/test/Dialect/Vector/invalid.mlir
index dcc4c75c72595..bc725b32be190 100644
--- a/mlir/test/Dialect/Vector/invalid.mlir
+++ b/mlir/test/Dialect/Vector/invalid.mlir
@@ -1639,6 +1639,20 @@ func.func @compress_memref_mismatch(%base: memref<?x?xf32>, %mask: vector<16xi1>
// -----
+func.func @compress_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
+ // expected-error @below {{'vector.compressstore' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ vector.compressstore %base[%c0], %mask, %value { alignment = -1 } : memref<?xf32>, vector<16xi1>, vector<16xf32>
+}
+
+// -----
+
+func.func @compress_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
+ // expected-error @below {{'vector.compressstore' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ vector.compressstore %base[%c0], %mask, %value { alignment = 3 } : memref<?xf32>, vector<16xi1>, vector<16xf32>
+}
+
+// -----
+
func.func @scan_reduction_dim_constraint(%arg0: vector<2x3xi32>, %arg1: vector<3xi32>) -> vector<3xi32> {
// expected-error at +1 {{'vector.scan' op reduction dimension 5 has to be less than 2}}
%0:2 = vector.scan <add>, %arg0, %arg1 {inclusive = true, reduction_dim = 5} :
>From a4d820f28053e716252a5d3eb634827c656ffda6 Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Wed, 6 Aug 2025 20:57:52 -0700
Subject: [PATCH 04/10] [mlir][vector] Add alignment to expandload
---
.../mlir/Dialect/Vector/IR/VectorOps.td | 27 ++++++++++++++++++-
mlir/test/Dialect/Vector/invalid.mlir | 14 ++++++++++
2 files changed, 40 insertions(+), 1 deletion(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index 0a36a54562fff..49cf4159b1268 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -2226,7 +2226,9 @@ def Vector_ExpandLoadOp :
Arguments<(ins Arg<AnyMemRef, "", [MemRead]>:$base,
Variadic<Index>:$indices,
FixedVectorOfNonZeroRankOf<[I1]>:$mask,
- AnyVectorOfNonZeroRank:$pass_thru)>,
+ AnyVectorOfNonZeroRank:$pass_thru,
+ ConfinedAttr<OptionalAttr<I64Attr>,
+ [AllAttrOf<[IntPositive, IntPowerOf2]>]>:$alignment)>,
Results<(outs AnyVectorOfNonZeroRank:$result)> {
let summary = "reads elements from memory and spreads them into a vector as defined by a mask";
@@ -2288,6 +2290,29 @@ def Vector_ExpandLoadOp :
"type($base) `,` type($mask) `,` type($pass_thru) `into` type($result)";
let hasCanonicalizer = 1;
let hasVerifier = 1;
+
+ let builders = [
+ OpBuilder<(ins "VectorType":$resultType,
+ "Value":$base,
+ "ValueRange":$indices,
+ "Value":$mask,
+ "Value":$passthrough,
+ CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ return build($_builder, $_state, resultType, base, indices, mask, passthrough,
+ alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ nullptr);
+ }]>,
+ OpBuilder<(ins "TypeRange":$resultTypes,
+ "Value":$base,
+ "ValueRange":$indices,
+ "Value":$mask,
+ "Value":$passthrough,
+ CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ return build($_builder, $_state, resultTypes, base, indices, mask, passthrough,
+ alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ nullptr);
+ }]>
+ ];
}
def Vector_CompressStoreOp :
diff --git a/mlir/test/Dialect/Vector/invalid.mlir b/mlir/test/Dialect/Vector/invalid.mlir
index bc725b32be190..68b07ec82aeb7 100644
--- a/mlir/test/Dialect/Vector/invalid.mlir
+++ b/mlir/test/Dialect/Vector/invalid.mlir
@@ -1607,6 +1607,20 @@ func.func @expand_memref_mismatch(%base: memref<?x?xf32>, %mask: vector<16xi1>,
// -----
+func.func @expand_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) {
+ // expected-error at +1 {{'vector.expandload' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ %0 = vector.expandload %base[%c0], %mask, %pass_thru { alignment = -1 } : memref<?xf32>, vector<16xi1>, vector<16xf32> into vector<16xf32>
+}
+
+// -----
+
+func.func @expand_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) {
+ // expected-error at +1 {{'vector.expandload' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
+ %0 = vector.expandload %base[%c0], %mask, %pass_thru { alignment = 3 } : memref<?xf32>, vector<16xi1>, vector<16xf32> into vector<16xf32>
+}
+
+// -----
+
func.func @compress_base_type_mismatch(%base: memref<?xf64>, %mask: vector<16xi1>, %value: vector<16xf32>) {
%c0 = arith.constant 0 : index
// expected-error at +1 {{'vector.compressstore' op base and valueToStore element type should match}}
>From e2ad0f90be5d7f11b03c697613f2544c9a8d9911 Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Thu, 7 Aug 2025 19:41:53 -0700
Subject: [PATCH 05/10] Use llvm::MaybeAlign
---
.../mlir/Dialect/Vector/IR/VectorOps.td | 24 +++++++++----------
1 file changed, 12 insertions(+), 12 deletions(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index 49cf4159b1268..aae2051600251 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -2121,9 +2121,9 @@ def Vector_GatherOp :
"Value":$index_vec,
"Value":$mask,
"Value":$passthrough,
- CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">:$alignment), [{
return build($_builder, $_state, resultType, base, indices, index_vec, mask, passthrough,
- alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
nullptr);
}]>,
OpBuilder<(ins "TypeRange":$resultTypes,
@@ -2132,9 +2132,9 @@ def Vector_GatherOp :
"Value":$index_vec,
"Value":$mask,
"Value":$passthrough,
- CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">:$alignment), [{
return build($_builder, $_state, resultTypes, base, indices, index_vec, mask, passthrough,
- alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
nullptr);
}]>
];
@@ -2213,9 +2213,9 @@ def Vector_ScatterOp :
"Value":$index_vec,
"Value":$mask,
"Value":$valueToStore,
- CArg<"llvm::Align", "llvm::Align()">: $alignment), [{
+ CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">: $alignment), [{
return build($_builder, $_state, base, indices, index_vec, mask, valueToStore,
- alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
nullptr);
}]>
];
@@ -2297,9 +2297,9 @@ def Vector_ExpandLoadOp :
"ValueRange":$indices,
"Value":$mask,
"Value":$passthrough,
- CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">:$alignment), [{
return build($_builder, $_state, resultType, base, indices, mask, passthrough,
- alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
nullptr);
}]>,
OpBuilder<(ins "TypeRange":$resultTypes,
@@ -2307,9 +2307,9 @@ def Vector_ExpandLoadOp :
"ValueRange":$indices,
"Value":$mask,
"Value":$passthrough,
- CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">:$alignment), [{
return build($_builder, $_state, resultTypes, base, indices, mask, passthrough,
- alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
nullptr);
}]>
];
@@ -2386,9 +2386,9 @@ def Vector_CompressStoreOp :
"ValueRange":$indices,
"Value":$mask,
"Value":$valueToStore,
- CArg<"llvm::Align", "llvm::Align()">:$alignment), [{
+ CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">:$alignment), [{
return build($_builder, $_state, base, indices, valueToStore, mask,
- alignment != llvm::Align() ? $_builder.getI64IntegerAttr(alignment.value()) :
+ alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
nullptr);
}]>
];
>From 92b3886401c724325cc6c4073ab8fe4ae340d48b Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Fri, 8 Aug 2025 20:42:54 -0700
Subject: [PATCH 06/10] Add documentation for alignment attribute
---
.../mlir/Dialect/Vector/IR/VectorOps.td | 24 +++++++++++++++----
1 file changed, 20 insertions(+), 4 deletions(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index aae2051600251..96c081ca2ad24 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -1382,6 +1382,11 @@ def Vector_TransferReadOp :
An additional `1` broadcast is required. On a GPU this broadcast could be
implemented using a warp-shuffle if loop `j` were mapped to `threadIdx.x`.
+ An optional `alignment` attribute allows to specify the byte alignment of the
+ gather operation. It must be a positive power of 2. The operation must access
+ memory at an address aligned to this boundary. Violations may lead to
+ architecture-specific faults or performance penalties.
+
Syntax
```
operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
@@ -1714,7 +1719,6 @@ def Vector_LoadOp : Vector_Op<"load", [
load operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
- A value of 0 indicates no specific alignment requirement.
}];
let arguments = (ins Arg<AnyMemRef, "the reference to load from",
@@ -1830,7 +1834,6 @@ def Vector_StoreOp : Vector_Op<"store", [
store operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
- A value of 0 indicates no specific alignment requirement.
}];
let arguments = (ins
@@ -1919,7 +1922,6 @@ def Vector_MaskedLoadOp :
load operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
- A value of 0 indicates no specific alignment requirement.
}];
let extraClassDeclaration = [{
MemRefType getMemRefType() {
@@ -2012,7 +2014,6 @@ def Vector_MaskedStoreOp :
store operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
- A value of 0 indicates no specific alignment requirement.
}];
let extraClassDeclaration = [{
MemRefType getMemRefType() {
@@ -2182,6 +2183,11 @@ def Vector_ScatterOp :
correspond to those of the `llvm.masked.scatter`
[intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-scatter-intrinsics).
+ An optional `alignment` attribute allows to specify the byte alignment of the
+ scatter operation. It must be a positive power of 2. The operation must access
+ memory at an address aligned to this boundary. Violations may lead to
+ architecture-specific faults or performance penalties.
+
Examples:
```mlir
@@ -2260,6 +2266,11 @@ def Vector_ExpandLoadOp :
correspond to those of the `llvm.masked.expandload`
[intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-expandload-intrinsics).
+ An optional `alignment` attribute allows to specify the byte alignment of the
+ load operation. It must be a positive power of 2. The operation must access
+ memory at an address aligned to this boundary. Violations may lead to
+ architecture-specific faults or performance penalties.
+
Note, at the moment this Op is only available for fixed-width vectors.
Examples:
@@ -2353,6 +2364,11 @@ def Vector_CompressStoreOp :
correspond to those of the `llvm.masked.compressstore`
[intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-compressstore-intrinsics).
+ An optional `alignment` attribute allows to specify the byte alignment of the
+ store operation. It must be a positive power of 2. The operation must access
+ memory at an address aligned to this boundary. Violations may lead to
+ architecture-specific faults or performance penalties.
+
Note, at the moment this Op is only available for fixed-width vectors.
Examples:
>From 482ad75b1d6e2f68fb045b22a283a9492c441e23 Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Mon, 11 Aug 2025 14:07:18 -0700
Subject: [PATCH 07/10] Change names in test functions
---
mlir/test/Dialect/Vector/invalid.mlir | 28 +++++++++++++--------------
1 file changed, 14 insertions(+), 14 deletions(-)
diff --git a/mlir/test/Dialect/Vector/invalid.mlir b/mlir/test/Dialect/Vector/invalid.mlir
index 68b07ec82aeb7..2e72bf036fa71 100644
--- a/mlir/test/Dialect/Vector/invalid.mlir
+++ b/mlir/test/Dialect/Vector/invalid.mlir
@@ -1317,7 +1317,7 @@ func.func @maskedload_negative_alignment(%base: memref<4xi32>, %mask: vector<32x
// -----
-func.func @maskedload_nonpoweroftwo_alignment(%base: memref<4xi32>, %mask: vector<32xi1>, %pass: vector<1xi32>, %index: index) {
+func.func @maskedload_non_power_of_2_alignment(%base: memref<4xi32>, %mask: vector<32xi1>, %pass: vector<1xi32>, %index: index) {
// expected-error at below {{'vector.maskedload' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
%val = vector.maskedload %base[%index], %mask, %pass { alignment = 3 } : memref<4xi32>, vector<32xi1>, vector<1xi32> into vector<1xi32>
return
@@ -1368,7 +1368,7 @@ func.func @maskedstore_negative_alignment(%base: memref<4xi32>, %mask: vector<32
// -----
-func.func @maskedstore_nonpoweroftwo_alignment(%base: memref<4xi32>, %mask: vector<32xi1>, %value: vector<1xi32>, %index: index) {
+func.func @maskedstore_non_power_of_2_alignment(%base: memref<4xi32>, %mask: vector<32xi1>, %value: vector<1xi32>, %index: index) {
// expected-error at below {{'vector.maskedstore' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
vector.maskedstore %base[%index], %mask, %value { alignment = 3 } : memref<4xi32>, vector<32xi1>, vector<1xi32> into vector<1xi32>
return
@@ -1470,7 +1470,7 @@ func.func @gather_pass_thru_type_mismatch(%base: memref<?xf32>, %indices: vector
// -----
-func.func @gather_invalid_alignment(%base: memref<16xf32>, %indices: vector<16xi32>,
+func.func @gather_negative_alignment(%base: memref<16xf32>, %indices: vector<16xi32>,
%mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0 : index) {
// expected-error at +2 {{'vector.gather' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
%0 = vector.gather %base[%c0][%indices], %mask, %pass_thru
@@ -1479,7 +1479,7 @@ func.func @gather_invalid_alignment(%base: memref<16xf32>, %indices: vector<16xi
// -----
-func.func @gather_invalid_alignment(%base: memref<16xf32>, %indices: vector<16xi32>,
+func.func @gather_non_power_of_two_alignment(%base: memref<16xf32>, %indices: vector<16xi32>,
%mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0 : index) {
// expected-error at +2 {{'vector.gather' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
%0 = vector.gather %base[%c0][%indices], %mask, %pass_thru
@@ -1549,7 +1549,7 @@ func.func @scatter_dim_mask_mismatch(%base: memref<?xf32>, %indices: vector<16xi
// -----
-func.func @scatter_invalid_alignment(%base: memref<?xf32>, %indices: vector<16xi32>,
+func.func @scatter_negative_alignment(%base: memref<?xf32>, %indices: vector<16xi32>,
%mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
// expected-error at +1 {{'vector.scatter' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
vector.scatter %base[%c0][%indices], %mask, %value { alignment = -1 }
@@ -1558,7 +1558,7 @@ func.func @scatter_invalid_alignment(%base: memref<?xf32>, %indices: vector<16xi
// -----
-func.func @scatter_invalid_alignment(%base: memref<?xf32>, %indices: vector<16xi32>,
+func.func @scatter_non_power_of_2_alignment(%base: memref<?xf32>, %indices: vector<16xi32>,
%mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
// expected-error at +1 {{'vector.scatter' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
vector.scatter %base[%c0][%indices], %mask, %value { alignment = 3 }
@@ -1607,14 +1607,14 @@ func.func @expand_memref_mismatch(%base: memref<?x?xf32>, %mask: vector<16xi1>,
// -----
-func.func @expand_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) {
+func.func @expand_negative_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) {
// expected-error at +1 {{'vector.expandload' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
%0 = vector.expandload %base[%c0], %mask, %pass_thru { alignment = -1 } : memref<?xf32>, vector<16xi1>, vector<16xf32> into vector<16xf32>
}
// -----
-func.func @expand_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) {
+func.func @expand_non_power_of_2_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) {
// expected-error at +1 {{'vector.expandload' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
%0 = vector.expandload %base[%c0], %mask, %pass_thru { alignment = 3 } : memref<?xf32>, vector<16xi1>, vector<16xf32> into vector<16xf32>
}
@@ -1653,14 +1653,14 @@ func.func @compress_memref_mismatch(%base: memref<?x?xf32>, %mask: vector<16xi1>
// -----
-func.func @compress_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
+func.func @compress_negative_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
// expected-error @below {{'vector.compressstore' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
vector.compressstore %base[%c0], %mask, %value { alignment = -1 } : memref<?xf32>, vector<16xi1>, vector<16xf32>
}
// -----
-func.func @compress_invalid_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
+func.func @compress_non_power_of_2_alignment(%base: memref<?xf32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) {
// expected-error @below {{'vector.compressstore' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
vector.compressstore %base[%c0], %mask, %value { alignment = 3 } : memref<?xf32>, vector<16xi1>, vector<16xf32>
}
@@ -2016,7 +2016,7 @@ func.func @vector_load(%src : memref<?xi8>) {
// -----
-func.func @invalid_load_alignment(%memref: memref<4xi32>, %c0: index) {
+func.func @load_negative_alignment(%memref: memref<4xi32>, %c0: index) {
// expected-error @below {{'vector.load' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
%val = vector.load %memref[%c0] { alignment = -1 } : memref<4xi32>, vector<4xi32>
return
@@ -2024,7 +2024,7 @@ func.func @invalid_load_alignment(%memref: memref<4xi32>, %c0: index) {
// -----
-func.func @invalid_load_alignment(%memref: memref<4xi32>, %c0: index) {
+func.func @load_non_pow_of_2_alignment(%memref: memref<4xi32>, %c0: index) {
// expected-error @below {{'vector.load' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
%val = vector.load %memref[%c0] { alignment = 3 } : memref<4xi32>, vector<4xi32>
return
@@ -2045,7 +2045,7 @@ func.func @vector_store(%dest : memref<?xi8>, %vec : vector<16x16xi8>) {
// -----
-func.func @invalid_store_alignment(%memref: memref<4xi32>, %val: vector<4xi32>, %c0: index) {
+func.func @store_negative_alignment(%memref: memref<4xi32>, %val: vector<4xi32>, %c0: index) {
// expected-error @below {{'vector.store' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
vector.store %val, %memref[%c0] { alignment = -1 } : memref<4xi32>, vector<4xi32>
return
@@ -2053,7 +2053,7 @@ func.func @invalid_store_alignment(%memref: memref<4xi32>, %val: vector<4xi32>,
// -----
-func.func @invalid_store_alignment(%memref: memref<4xi32>, %val: vector<4xi32>, %c0: index) {
+func.func @store_non_pow_of_2_alignment(%memref: memref<4xi32>, %val: vector<4xi32>, %c0: index) {
// expected-error @below {{'vector.store' op attribute 'alignment' failed to satisfy constraint: 64-bit signless integer attribute whose value is positive and whose value is a power of two > 0}}
vector.store %val, %memref[%c0] { alignment = 3 } : memref<4xi32>, vector<4xi32>
return
>From 9b451db43c23ab4a8441790ba05d1dccff0cf84a Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Mon, 11 Aug 2025 14:24:59 -0700
Subject: [PATCH 08/10] Remove unnecessary constructor
---
mlir/include/mlir/Dialect/Vector/IR/VectorOps.td | 10 ----------
1 file changed, 10 deletions(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index 96c081ca2ad24..22be7f25379a3 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -2312,16 +2312,6 @@ def Vector_ExpandLoadOp :
return build($_builder, $_state, resultType, base, indices, mask, passthrough,
alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
nullptr);
- }]>,
- OpBuilder<(ins "TypeRange":$resultTypes,
- "Value":$base,
- "ValueRange":$indices,
- "Value":$mask,
- "Value":$passthrough,
- CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">:$alignment), [{
- return build($_builder, $_state, resultTypes, base, indices, mask, passthrough,
- alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) :
- nullptr);
}]>
];
}
>From 3180cd0c1875bb3c3e2b8ebd561ea0a9acb7ba1d Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Mon, 11 Aug 2025 14:26:18 -0700
Subject: [PATCH 09/10] Remove documentation from wrong op
---
mlir/include/mlir/Dialect/Vector/IR/VectorOps.td | 5 -----
1 file changed, 5 deletions(-)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index 22be7f25379a3..79a216a853a16 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -1382,11 +1382,6 @@ def Vector_TransferReadOp :
An additional `1` broadcast is required. On a GPU this broadcast could be
implemented using a warp-shuffle if loop `j` were mapped to `threadIdx.x`.
- An optional `alignment` attribute allows to specify the byte alignment of the
- gather operation. It must be a positive power of 2. The operation must access
- memory at an address aligned to this boundary. Violations may lead to
- architecture-specific faults or performance penalties.
-
Syntax
```
operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
>From 47db5b173ffaddbc7a723d79881017a9ae2818a0 Mon Sep 17 00:00:00 2001
From: Erick Ochoa <erick.ochoalopez at amd.com>
Date: Mon, 11 Aug 2025 14:32:31 -0700
Subject: [PATCH 10/10] Fix documention
---
mlir/include/mlir/Dialect/Vector/IR/VectorOps.td | 13 +++++++++++++
1 file changed, 13 insertions(+)
diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
index 79a216a853a16..ddcb00cffde63 100644
--- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
+++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td
@@ -1714,6 +1714,7 @@ def Vector_LoadOp : Vector_Op<"load", [
load operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
}];
let arguments = (ins Arg<AnyMemRef, "the reference to load from",
@@ -1829,6 +1830,7 @@ def Vector_StoreOp : Vector_Op<"store", [
store operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
}];
let arguments = (ins
@@ -1917,6 +1919,7 @@ def Vector_MaskedLoadOp :
load operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
}];
let extraClassDeclaration = [{
MemRefType getMemRefType() {
@@ -2009,6 +2012,7 @@ def Vector_MaskedStoreOp :
store operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
}];
let extraClassDeclaration = [{
MemRefType getMemRefType() {
@@ -2083,6 +2087,12 @@ def Vector_GatherOp :
during progressively lowering to bring other memory operations closer to
hardware ISA support for a gather.
+ An optional `alignment` attribute allows to specify the byte alignment of the
+ scatter operation. It must be a positive power of 2. The operation must access
+ memory at an address aligned to this boundary. Violations may lead to
+ architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
+
Examples:
```mlir
@@ -2182,6 +2192,7 @@ def Vector_ScatterOp :
scatter operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
Examples:
@@ -2265,6 +2276,7 @@ def Vector_ExpandLoadOp :
load operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
Note, at the moment this Op is only available for fixed-width vectors.
@@ -2353,6 +2365,7 @@ def Vector_CompressStoreOp :
store operation. It must be a positive power of 2. The operation must access
memory at an address aligned to this boundary. Violations may lead to
architecture-specific faults or performance penalties.
+ A value of 0 indicates no specific alignment requirement.
Note, at the moment this Op is only available for fixed-width vectors.
More information about the Mlir-commits
mailing list