From a97cbc6a1223fd34dc9e12e6fe8f4d3445fac47e Mon Sep 17 00:00:00 2001 From: Erick Ochoa Date: Wed, 6 Aug 2025 19:48:16 -0700 Subject: [PATCH 01/11] [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:$indices, VectorOfNonZeroRankOf<[AnyInteger, Index]>:$index_vec, VectorOfNonZeroRankOf<[I1]>:$mask, - AnyVectorOfNonZeroRank:$pass_thru)>, + AnyVectorOfNonZeroRank:$pass_thru, + ConfinedAttr, + [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, %indices: vector // ----- +func.func @gather_invalid_alignment(%base: memref<16xf32>, %indices: vector<16xi32>, + %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0 : index) { + // expected-error@+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@+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 Date: Wed, 6 Aug 2025 19:58:43 -0700 Subject: [PATCH 02/11] [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:$indices, VectorOfNonZeroRankOf<[AnyInteger, Index]>:$index_vec, VectorOfNonZeroRankOf<[I1]>:$mask, - AnyVectorOfNonZeroRank:$valueToStore)> { + AnyVectorOfNonZeroRank:$valueToStore, + ConfinedAttr, + [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, %indices: vector<16xi // ----- +func.func @scatter_invalid_alignment(%base: memref, %indices: vector<16xi32>, + %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) { + // expected-error@+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, vector<16xi32>, vector<16xi1>, vector<16xf32> +} + +// ----- + +func.func @scatter_invalid_alignment(%base: memref, %indices: vector<16xi32>, + %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) { + // expected-error@+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, vector<16xi32>, vector<16xi1>, vector<16xf32> +} + +// ----- + func.func @expand_base_type_mismatch(%base: memref, %mask: vector<16xi1>, %pass_thru: vector<16xf32>) { %c0 = arith.constant 0 : index // expected-error@+1 {{'vector.expandload' op base and result element type should match}} From 36949d1cbce78fc518619c759c557afd1feccd3f Mon Sep 17 00:00:00 2001 From: Erick Ochoa Date: Wed, 6 Aug 2025 20:48:21 -0700 Subject: [PATCH 03/11] [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:$base, Variadic:$indices, FixedVectorOfNonZeroRankOf<[I1]>:$mask, - AnyVectorOfNonZeroRank:$valueToStore)> { + AnyVectorOfNonZeroRank:$valueToStore, + ConfinedAttr, + [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, %mask: vector<16xi1> // ----- +func.func @compress_invalid_alignment(%base: memref, %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, vector<16xi1>, vector<16xf32> +} + +// ----- + +func.func @compress_invalid_alignment(%base: memref, %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, vector<16xi1>, vector<16xf32> +} + +// ----- + func.func @scan_reduction_dim_constraint(%arg0: vector<2x3xi32>, %arg1: vector<3xi32>) -> vector<3xi32> { // expected-error@+1 {{'vector.scan' op reduction dimension 5 has to be less than 2}} %0:2 = vector.scan , %arg0, %arg1 {inclusive = true, reduction_dim = 5} : From a4d820f28053e716252a5d3eb634827c656ffda6 Mon Sep 17 00:00:00 2001 From: Erick Ochoa Date: Wed, 6 Aug 2025 20:57:52 -0700 Subject: [PATCH 04/11] [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:$base, Variadic:$indices, FixedVectorOfNonZeroRankOf<[I1]>:$mask, - AnyVectorOfNonZeroRank:$pass_thru)>, + AnyVectorOfNonZeroRank:$pass_thru, + ConfinedAttr, + [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, %mask: vector<16xi1>, // ----- +func.func @expand_invalid_alignment(%base: memref, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) { + // expected-error@+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, vector<16xi1>, vector<16xf32> into vector<16xf32> +} + +// ----- + +func.func @expand_invalid_alignment(%base: memref, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) { + // expected-error@+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, vector<16xi1>, vector<16xf32> into vector<16xf32> +} + +// ----- + func.func @compress_base_type_mismatch(%base: memref, %mask: vector<16xi1>, %value: vector<16xf32>) { %c0 = arith.constant 0 : index // expected-error@+1 {{'vector.compressstore' op base and valueToStore element type should match}} From e2ad0f90be5d7f11b03c697613f2544c9a8d9911 Mon Sep 17 00:00:00 2001 From: Erick Ochoa Date: Thu, 7 Aug 2025 19:41:53 -0700 Subject: [PATCH 05/11] 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 Date: Fri, 8 Aug 2025 20:42:54 -0700 Subject: [PATCH 06/11] 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 Date: Mon, 11 Aug 2025 14:07:18 -0700 Subject: [PATCH 07/11] 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@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@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, %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@+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@+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, %indices: vector<16xi // ----- -func.func @scatter_invalid_alignment(%base: memref, %indices: vector<16xi32>, +func.func @scatter_negative_alignment(%base: memref, %indices: vector<16xi32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) { // expected-error@+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, %indices: vector<16xi // ----- -func.func @scatter_invalid_alignment(%base: memref, %indices: vector<16xi32>, +func.func @scatter_non_power_of_2_alignment(%base: memref, %indices: vector<16xi32>, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) { // expected-error@+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, %mask: vector<16xi1>, // ----- -func.func @expand_invalid_alignment(%base: memref, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) { +func.func @expand_negative_alignment(%base: memref, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) { // expected-error@+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, vector<16xi1>, vector<16xf32> into vector<16xf32> } // ----- -func.func @expand_invalid_alignment(%base: memref, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) { +func.func @expand_non_power_of_2_alignment(%base: memref, %mask: vector<16xi1>, %pass_thru: vector<16xf32>, %c0: index) { // expected-error@+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, vector<16xi1>, vector<16xf32> into vector<16xf32> } @@ -1653,14 +1653,14 @@ func.func @compress_memref_mismatch(%base: memref, %mask: vector<16xi1> // ----- -func.func @compress_invalid_alignment(%base: memref, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) { +func.func @compress_negative_alignment(%base: memref, %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, vector<16xi1>, vector<16xf32> } // ----- -func.func @compress_invalid_alignment(%base: memref, %mask: vector<16xi1>, %value: vector<16xf32>, %c0: index) { +func.func @compress_non_power_of_2_alignment(%base: memref, %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, vector<16xi1>, vector<16xf32> } @@ -2016,7 +2016,7 @@ func.func @vector_load(%src : memref) { // ----- -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, %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 Date: Mon, 11 Aug 2025 14:24:59 -0700 Subject: [PATCH 08/11] 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 Date: Mon, 11 Aug 2025 14:26:18 -0700 Subject: [PATCH 09/11] 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 Date: Mon, 11 Aug 2025 14:32:31 -0700 Subject: [PATCH 10/11] 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 Date: Mon, 11 Aug 2025 14:38:25 -0700 Subject: [PATCH 11/11] Remove unnecessary constructor --- mlir/include/mlir/Dialect/Vector/IR/VectorOps.td | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td index ddcb00cffde63..78def9cc78520 100644 --- a/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td +++ b/mlir/include/mlir/Dialect/Vector/IR/VectorOps.td @@ -2131,17 +2131,6 @@ def Vector_GatherOp : return build($_builder, $_state, resultType, base, indices, index_vec, mask, passthrough, alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) : nullptr); - }]>, - OpBuilder<(ins "TypeRange":$resultTypes, - "Value":$base, - "ValueRange":$indices, - "Value":$index_vec, - "Value":$mask, - "Value":$passthrough, - CArg<"llvm::MaybeAlign", "llvm::MaybeAlign()">:$alignment), [{ - return build($_builder, $_state, resultTypes, base, indices, index_vec, mask, passthrough, - alignment.has_value() ? $_builder.getI64IntegerAttr(alignment->value()) : - nullptr); }]> ]; }