1//===- VectorOps.td - Vector op definitions ---------------*- tablegen -*-====// 2// 3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4// See https://llvm.org/LICENSE.txt for license information. 5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6// 7//===----------------------------------------------------------------------===// 8// 9// Defines MLIR vector operations. 10// 11//===----------------------------------------------------------------------===// 12 13#ifndef VECTOR_OPS 14#define VECTOR_OPS 15 16include "mlir/IR/EnumAttr.td" 17include "mlir/Interfaces/ControlFlowInterfaces.td" 18include "mlir/Interfaces/InferTypeOpInterface.td" 19include "mlir/Interfaces/SideEffectInterfaces.td" 20include "mlir/Interfaces/VectorInterfaces.td" 21include "mlir/Interfaces/ViewLikeInterface.td" 22 23def Vector_Dialect : Dialect { 24 let name = "vector"; 25 let cppNamespace = "::mlir::vector"; 26 27 let useDefaultAttributePrinterParser = 1; 28 let hasConstantMaterializer = 1; 29 let dependentDialects = ["arith::ArithmeticDialect"]; 30 let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed; 31} 32 33// Base class for Vector dialect ops. 34class Vector_Op<string mnemonic, list<Trait> traits = []> : 35 Op<Vector_Dialect, mnemonic, traits>; 36 37// The "kind" of combining function for contractions and reductions. 38def COMBINING_KIND_ADD : I32BitEnumAttrCaseBit<"ADD", 0, "add">; 39def COMBINING_KIND_MUL : I32BitEnumAttrCaseBit<"MUL", 1, "mul">; 40def COMBINING_KIND_MINUI : I32BitEnumAttrCaseBit<"MINUI", 2, "minui">; 41def COMBINING_KIND_MINSI : I32BitEnumAttrCaseBit<"MINSI", 3, "minsi">; 42def COMBINING_KIND_MINF : I32BitEnumAttrCaseBit<"MINF", 4, "minf">; 43def COMBINING_KIND_MAXUI : I32BitEnumAttrCaseBit<"MAXUI", 5, "maxui">; 44def COMBINING_KIND_MAXSI : I32BitEnumAttrCaseBit<"MAXSI", 6, "maxsi">; 45def COMBINING_KIND_MAXF : I32BitEnumAttrCaseBit<"MAXF", 7, "maxf">; 46def COMBINING_KIND_AND : I32BitEnumAttrCaseBit<"AND", 8, "and">; 47def COMBINING_KIND_OR : I32BitEnumAttrCaseBit<"OR", 9, "or">; 48def COMBINING_KIND_XOR : I32BitEnumAttrCaseBit<"XOR", 10, "xor">; 49 50def CombiningKind : I32BitEnumAttr< 51 "CombiningKind", 52 "Kind of combining function for contractions and reductions", 53 [COMBINING_KIND_ADD, COMBINING_KIND_MUL, COMBINING_KIND_MINUI, 54 COMBINING_KIND_MINSI, COMBINING_KIND_MINF, COMBINING_KIND_MAXUI, 55 COMBINING_KIND_MAXSI, COMBINING_KIND_MAXF, COMBINING_KIND_AND, 56 COMBINING_KIND_OR, COMBINING_KIND_XOR]> { 57 let cppNamespace = "::mlir::vector"; 58 let genSpecializedAttr = 0; 59} 60 61def Vector_CombiningKindAttr : DialectAttr< 62 Vector_Dialect, 63 CPred<"$_self.isa<::mlir::vector::CombiningKindAttr>()">, 64 "Kind of combining function for contractions and reductions"> { 65 let storageType = "::mlir::vector::CombiningKindAttr"; 66 let returnType = "::mlir::vector::CombiningKind"; 67 let convertFromStorage = "$_self.getKind()"; 68 let constBuilderCall = 69 "::mlir::vector::CombiningKindAttr::get($0, $_builder.getContext())"; 70} 71 72// TODO: Add an attribute to specify a different algebra with operators other 73// than the current set: {*, +}. 74def Vector_ContractionOp : 75 Vector_Op<"contract", [ 76 NoSideEffect, 77 PredOpTrait<"lhs and rhs have same element type", TCopVTEtIsSameAs<0, 1>>, 78 PredOpTrait<"third operand acc and result have same element type", 79 TCresVTEtIsSameAsOpBase<0, 2>>, 80 DeclareOpInterfaceMethods<VectorUnrollOpInterface, ["getShapeForUnroll"]> 81 ]>, 82 Arguments<(ins AnyVector:$lhs, AnyVector:$rhs, AnyType:$acc, 83 Variadic<VectorOf<[I1]>>:$masks, 84 ArrayAttr:$indexing_maps, 85 ArrayAttr:$iterator_types, 86 DefaultValuedAttr<Vector_CombiningKindAttr, 87 "CombiningKind::ADD">:$kind)>, 88 Results<(outs AnyType)> { 89 let summary = "vector contraction operation"; 90 let description = [{ 91 Computes the sum of products of vector elements along contracting 92 dimension pairs from 2 vectors of rank M and N respectively, adds this 93 intermediate result to the accumulator argument of rank K, and returns a 94 vector result of rank K (where K = num_lhs_free_dims + num_rhs_free_dims + 95 num_batch_dims (see dimension type descriptions below)). For K = 0 (no 96 free or batch dimensions), the accumulator and output are a scalar. 97 98 Optional vector mask arguments (produced by CreateMaskOp or ConstantMaskOp) 99 specify the dynamic dimension sizes of valid data within the lhs/rhs vector 100 arguments. 101 102 An iterator type attribute list must be specified, where each element of 103 the list represents an iterator with one of the following types: 104 105 *) "reduction": reduction dimensions are present in the lhs and rhs 106 arguments but not in the output (and accumulator 107 argument). These are the dimensions along which the vector 108 contraction op computes the sum of products, and 109 contracting dimension pair dimension sizes must match 110 between lhs/rhs. 111 *) "parallel": Batch dimensions are iterator type "parallel", and 112 are non-contracting dimensions present in the lhs, rhs and 113 output. The lhs/rhs co-iterate along the batch dimensions, 114 which should be expressed in their indexing maps. 115 116 Free dimensions are iterator type "parallel", and are 117 non-contraction, non-batch dimensions accessed by either the 118 lhs or rhs (but not both). The lhs and rhs free dimensions 119 are unrelated to each other and do not co-iterate, which 120 should be expressed in their indexing maps. 121 122 An indexing map attribute list must be specified with an entry for lhs, rhs 123 and acc arguments. An indexing map attribute specifies a mapping from each 124 iterator in the iterator type list, to each dimension of an N-D vector. 125 126 An optional kind attribute may be used to specify the combining function 127 between the intermediate result and accumulator argument of rank K. This 128 attribute can take the values add/mul/min/max for int/fp, and/or/xor for 129 int only. The default is "add". 130 131 Example: 132 133 ```mlir 134 // Simple DOT product (K = 0). 135 #contraction_accesses = [ 136 affine_map<(i) -> (i)>, 137 affine_map<(i) -> (i)>, 138 affine_map<(i) -> ()> 139 ] 140 #contraction_trait = { 141 indexing_maps = #contraction_accesses, 142 iterator_types = ["reduction"] 143 } 144 %3 = vector.contract #contraction_trait %0, %1, %2 145 : vector<10xf32>, vector<10xf32> into f32 146 147 // 2D vector contraction with one contracting dimension (matmul, K = 2). 148 #contraction_accesses = [ 149 affine_map<(i, j, k) -> (i, k)>, 150 affine_map<(i, j, k) -> (k, j)>, 151 affine_map<(i, j, k) -> (i, j)> 152 ] 153 #contraction_trait = { 154 indexing_maps = #contraction_accesses, 155 iterator_types = ["parallel", "parallel", "reduction"] 156 } 157 158 %3 = vector.contract #contraction_trait %0, %1, %2 159 : vector<4x3xf32>, vector<3x7xf32> into vector<4x7xf32> 160 161 // 4D to 3D vector contraction with two contracting dimensions and 162 // one batch dimension (K = 3). 163 #contraction_accesses = [ 164 affine_map<(b0, f0, f1, c0, c1) -> (c0, b0, c1, f0)>, 165 affine_map<(b0, f0, f1, c0, c1) -> (b0, c1, c0, f1)>, 166 affine_map<(b0, f0, f1, c0, c1) -> (b0, f0, f1)> 167 ] 168 #contraction_trait = { 169 indexing_maps = #contraction_accesses, 170 iterator_types = ["parallel", "parallel", "parallel", 171 "reduction", "reduction"] 172 } 173 174 %4 = vector.contract #contraction_trait %0, %1, %2 175 : vector<7x8x16x15xf32>, vector<8x16x7x5xf32> into vector<8x15x5xf32> 176 177 // 4D vector contraction with two contracting dimensions and optional 178 // vector mask arguments. 179 %lhs_mask = vector.constant_mask [7, 8, 16, 15] : vector<7x8x16x15xi1> 180 %rhs_mask = vector.constant_mask [8, 16, 7, 5] : vector<8x16x7x5xi1> 181 182 %5 = vector.contract #contraction_trait %0, %1, %2, %lhs_mask, %rhs_mask 183 : vector<7x8x16x15xf32>, vector<8x16x7x5xf32> into vector<8x15x8x5xf32> 184 185 // Vector contraction with mixed typed. lhs/rhs have different element 186 // types than accumulator/result. 187 %6 = vector.contract #contraction_trait %0, %1, %2 188 : vector<10xf16>, vector<10xf16> into f32 189 190 // Contract with max (K = 0). 191 #contraction_accesses = [ 192 affine_map<(i) -> (i)>, 193 affine_map<(i) -> (i)>, 194 affine_map<(i) -> ()> 195 ] 196 #contraction_trait = { 197 indexing_maps = #contraction_accesses, 198 iterator_types = ["reduction"], 199 kind = #vector.kind<max> 200 } 201 %7 = vector.contract #contraction_trait %0, %1, %2 202 : vector<10xf32>, vector<10xf32> into f32 203 ``` 204 }]; 205 let builders = [ 206 OpBuilder<(ins "Value":$lhs, "Value":$rhs, "Value":$acc, 207 "ArrayAttr":$indexingMaps, "ArrayAttr":$iteratorTypes)>, 208 OpBuilder<(ins "Value":$lhs, "Value":$rhs, "Value":$acc, 209 "ArrayRef<ArrayRef<AffineExpr>>":$indexingExprs, 210 "ArrayRef<StringRef>":$iteratorTypes)>, 211 OpBuilder<(ins "Value":$lhs, "Value":$rhs, "Value":$acc, 212 "ArrayAttr":$indexingMaps, "ArrayAttr":$iteratorTypes, 213 "CombiningKind":$kind)> 214 ]; 215 let extraClassDeclaration = [{ 216 VectorType getLhsType() { 217 return getLhs().getType().cast<VectorType>(); 218 } 219 VectorType getRhsType() { 220 return getRhs().getType().cast<VectorType>(); 221 } 222 Type getAccType() { return getAcc().getType(); } 223 VectorType getLHSVectorMaskType() { 224 if (llvm::size(getMasks()) != 2) return VectorType(); 225 return getOperand(3).getType().cast<VectorType>(); 226 } 227 VectorType getRHSVectorMaskType() { 228 if (llvm::size(getMasks()) != 2) return VectorType(); 229 return getOperand(4).getType().cast<VectorType>(); 230 } 231 Type getResultType() { return getResult().getType(); } 232 ArrayRef<StringRef> getTraitAttrNames(); 233 static unsigned getAccOperandIndex() { return 2; } 234 235 llvm::SmallVector<::mlir::AffineMap, 4> getIndexingMapsArray() { 236 return llvm::to_vector<4>(getIndexingMaps().getAsValueRange<::mlir::AffineMapAttr>()); 237 } 238 239 // Returns the bounds of each dimension in the iteration space spanned 240 // by the iterator types of this operation. 241 void getIterationBounds(SmallVectorImpl<int64_t> &iterationBounds); 242 243 // Returns a list of index maps, where there is a list entry for each 244 // op indexing map attribute (i.e. one for each input and output, with 245 // the output listed last). Each index map, maps from this operations 246 // iteration space, to vector dimensions of the maps input/output. 247 void getIterationIndexMap( 248 std::vector<DenseMap<int64_t, int64_t>> &iterationIndexMap); 249 250 std::vector<std::pair<int64_t, int64_t>> getContractingDimMap(); 251 std::vector<std::pair<int64_t, int64_t>> getBatchDimMap(); 252 253 static constexpr StringRef getKindAttrStrName() { return "kind"; } 254 255 static CombiningKind getDefaultKind() { 256 return CombiningKind::ADD; 257 } 258 }]; 259 260 let hasCanonicalizer = 1; 261 let hasCustomAssemblyFormat = 1; 262 let hasVerifier = 1; 263} 264 265def Vector_ReductionOp : 266 Vector_Op<"reduction", [NoSideEffect, 267 PredOpTrait<"source operand and result have same element type", 268 TCresVTEtIsSameAsOpBase<0, 0>>, 269 DeclareOpInterfaceMethods<VectorUnrollOpInterface, 270 ["getShapeForUnroll"]>]>, 271 Arguments<(ins Vector_CombiningKindAttr:$kind, AnyVector:$vector, 272 Optional<AnyType>:$acc)>, 273 Results<(outs AnyType:$dest)> { 274 let summary = "reduction operation"; 275 let description = [{ 276 Reduces an 1-D vector "horizontally" into a scalar using the given 277 operation (add/mul/min/max for int/fp and and/or/xor for int only). 278 Reductions also allow an optional fused accumulator. 279 280 Note that these operations are restricted to 1-D vectors to remain 281 close to the corresponding LLVM intrinsics: 282 283 http://llvm.org/docs/LangRef.html#vector-reduction-intrinsics 284 285 Example: 286 287 ```mlir 288 %1 = vector.reduction <add>, %0 : vector<16xf32> into f32 289 290 %3 = vector.reduction <xor>, %2 : vector<4xi32> into i32 291 292 %4 = vector.reduction <mul>, %0, %1 : vector<16xf32> into f32 293 ``` 294 }]; 295 let extraClassDeclaration = [{ 296 VectorType getVectorType() { 297 return getVector().getType().cast<VectorType>(); 298 } 299 }]; 300 let builders = [ 301 // Builder that infers the type of `dest`. 302 OpBuilder<(ins "CombiningKind":$kind, "Value":$vector, "Value":$acc)>, 303 // Builder that infers the type of `dest` and has no accumulator. 304 OpBuilder<(ins "CombiningKind":$kind, "Value":$vector)> 305 ]; 306 307 // TODO: Migrate to assemblyFormat once `AllTypesMatch` supports optional 308 // operands. 309 let hasCustomAssemblyFormat = 1; 310 let hasCanonicalizer = 1; 311 let hasVerifier = 1; 312} 313 314def Vector_MultiDimReductionOp : 315 Vector_Op<"multi_reduction", [NoSideEffect, 316 AllTypesMatch<["dest", "acc"]>, 317 PredOpTrait<"source operand and result have same element type", 318 TCresVTEtIsSameAsOpBase<0, 0>>, 319 DeclareOpInterfaceMethods<InferTypeOpInterface>, 320 DeclareOpInterfaceMethods<VectorUnrollOpInterface, 321 ["getShapeForUnroll"]>]>, 322 Arguments<(ins Vector_CombiningKindAttr:$kind, 323 AnyVector:$source, 324 AnyType:$acc, 325 I64ArrayAttr:$reduction_dims)>, 326 Results<(outs AnyType:$dest)> { 327 let summary = "Multi-dimensional reduction operation"; 328 let description = [{ 329 Reduces an n-D vector into an (n-k)-D vector (or a scalar when k == n) 330 using the given operation (add/mul/min/max for int/fp and and/or/xor for 331 int only). 332 Takes an initial accumulator operand. 333 334 Example: 335 336 ```mlir 337 %1 = vector.multi_reduction <add>, %0, %acc0 [1, 3] : 338 vector<4x8x16x32xf32> into vector<4x16xf32> 339 %2 = vector.multi_reduction <add>, %1, %acc1 [0, 1] : 340 vector<4x16xf32> into f32 341 ``` 342 }]; 343 let builders = [ 344 OpBuilder<(ins "Value":$source, "Value":$acc, 345 "ArrayRef<bool>":$reductionMask, "CombiningKind":$kind)> 346 ]; 347 let extraClassDeclaration = [{ 348 static StringRef getKindAttrStrName() { return "kind"; } 349 static StringRef getReductionDimsAttrStrName() { return "reduction_dims"; } 350 351 VectorType getSourceVectorType() { 352 return getSource().getType().cast<VectorType>(); 353 } 354 Type getDestType() { 355 return getDest().getType(); 356 } 357 358 bool isReducedDim(int64_t d) { 359 assert(d >= 0 && d < static_cast<int64_t>(getReductionMask().size()) && 360 "d overflows the number of dims"); 361 return getReductionMask()[d]; 362 } 363 364 SmallVector<bool> getReductionMask() { 365 SmallVector<bool> res(getSourceVectorType().getRank(), false); 366 for (auto ia : getReductionDims().getAsRange<IntegerAttr>()) 367 res[ia.getInt()] = true; 368 return res; 369 } 370 static SmallVector<bool> getReductionMask( 371 ArrayRef<int64_t> reductionDims, unsigned sourceRank) { 372 SmallVector<bool> res(sourceRank, false); 373 for (auto idx : reductionDims) 374 res[idx] = true; 375 return res; 376 } 377 }]; 378 let assemblyFormat = 379 "$kind `,` $source `,` $acc attr-dict $reduction_dims `:` type($source) `to` type($dest)"; 380 let hasFolder = 1; 381 let hasVerifier = 1; 382} 383 384def Vector_BroadcastOp : 385 Vector_Op<"broadcast", [NoSideEffect, 386 PredOpTrait<"source operand and result have same element type", 387 TCresVTEtIsSameAsOpBase<0, 0>>]>, 388 Arguments<(ins AnyType:$source)>, 389 Results<(outs AnyVectorOfAnyRank:$vector)> { 390 let summary = "broadcast operation"; 391 let description = [{ 392 Broadcasts the scalar or k-D vector value in the source operand 393 to a n-D result vector such that the broadcast makes sense, i.e., 394 the source operand is duplicated to match the given rank and sizes 395 in the result vector. The legality rules are: 396 * the source operand must have the same element type as the result type 397 * a k-D vector <s_1 x .. x s_k x type> can be broadcast to 398 a n-D vector <t_1 x .. x t_n x type> if 399 * k <= n, and 400 * the sizes in the trailing dimensions n-k < i <= n with j=i+k-n 401 match exactly as s_j = t_i or s_j = 1: 402 ``` 403 t_1 x .. t_n-k x t_n-k+1 x .. x t_i x .. x t_n 404 s_1 x .. x s_j x .. x s_k 405 <duplication> <potential stretch> 406 ``` 407 The source operand is duplicated over all the missing leading dimensions 408 and stretched over the trailing dimensions where the source has a non-equal 409 dimension of 1. These rules imply that any scalar broadcast (k=0) to any 410 shaped vector with the same element type is always legal. 411 412 Example: 413 414 ```mlir 415 %0 = arith.constant 0.0 : f32 416 %1 = vector.broadcast %0 : f32 to vector<16xf32> 417 %2 = vector.broadcast %1 : vector<16xf32> to vector<4x16xf32> 418 ``` 419 }]; 420 let extraClassDeclaration = [{ 421 Type getSourceType() { return getSource().getType(); } 422 VectorType getVectorType() { 423 return getVector().getType().cast<VectorType>(); 424 } 425 }]; 426 let assemblyFormat = "$source attr-dict `:` type($source) `to` type($vector)"; 427 let hasFolder = 1; 428 let hasCanonicalizer = 1; 429 let hasVerifier = 1; 430} 431 432def Vector_ShuffleOp : 433 Vector_Op<"shuffle", [NoSideEffect, 434 PredOpTrait<"first operand v1 and result have same element type", 435 TCresVTEtIsSameAsOpBase<0, 0>>, 436 PredOpTrait<"second operand v2 and result have same element type", 437 TCresVTEtIsSameAsOpBase<0, 1>>, 438 DeclareOpInterfaceMethods<InferTypeOpInterface>]>, 439 Arguments<(ins AnyVector:$v1, AnyVector:$v2, I64ArrayAttr:$mask)>, 440 Results<(outs AnyVector:$vector)> { 441 let summary = "shuffle operation"; 442 let description = [{ 443 The shuffle operation constructs a permutation (or duplication) of elements 444 from two input vectors, returning a vector with the same element type as 445 the input and a length that is the same as the shuffle mask. The two input 446 vectors must have the same element type, rank, and trailing dimension sizes 447 and shuffles their values in the leading dimension (which may differ in size) 448 according to the given mask. The legality rules are: 449 * the two operands must have the same element type as the result 450 * the two operands and the result must have the same rank and trailing 451 dimension sizes, viz. given two k-D operands 452 v1 : <s_1 x s_2 x .. x s_k x type> and 453 v2 : <t_1 x t_2 x .. x t_k x type> 454 we have s_i = t_i for all 1 < i <= k 455 * the mask length equals the leading dimension size of the result 456 * numbering the input vector indices left to right across the operands, all 457 mask values must be within range, viz. given two k-D operands v1 and v2 458 above, all mask values are in the range [0,s_1+t_1) 459 460 Example: 461 462 ```mlir 463 %0 = vector.shuffle %a, %b[0, 3] 464 : vector<2xf32>, vector<2xf32> ; yields vector<2xf32> 465 %1 = vector.shuffle %c, %b[0, 1, 2] 466 : vector<2x16xf32>, vector<1x16xf32> ; yields vector<3x16xf32> 467 %2 = vector.shuffle %a, %b[3, 2, 1, 0] 468 : vector<2xf32>, vector<2xf32> ; yields vector<4xf32> 469 ``` 470 }]; 471 let builders = [ 472 OpBuilder<(ins "Value":$v1, "Value":$v2, "ArrayRef<int64_t>")> 473 ]; 474 let hasFolder = 1; 475 let extraClassDeclaration = [{ 476 static StringRef getMaskAttrStrName() { return "mask"; } 477 VectorType getV1VectorType() { 478 return getV1().getType().cast<VectorType>(); 479 } 480 VectorType getV2VectorType() { 481 return getV2().getType().cast<VectorType>(); 482 } 483 VectorType getVectorType() { 484 return getVector().getType().cast<VectorType>(); 485 } 486 }]; 487 let assemblyFormat = "operands $mask attr-dict `:` type(operands)"; 488 let hasVerifier = 1; 489 let hasCanonicalizer = 1; 490} 491 492def Vector_ExtractElementOp : 493 Vector_Op<"extractelement", [NoSideEffect, 494 TypesMatchWith<"result type matches element type of vector operand", 495 "vector", "result", 496 "$_self.cast<ShapedType>().getElementType()">]>, 497 Arguments<(ins AnyVectorOfAnyRank:$vector, 498 Optional<AnySignlessIntegerOrIndex>:$position)>, 499 Results<(outs AnyType:$result)> { 500 let summary = "extractelement operation"; 501 let description = [{ 502 Takes a 0-D or 1-D vector and a optional dynamic index position and 503 extracts the scalar at that position. 504 505 Note that this instruction resembles vector.extract, but is restricted to 506 0-D and 1-D vectors and relaxed to dynamic indices. 507 If the vector is 0-D, the position must be llvm::None. 508 509 510 It is meant to be closer to LLVM's version: 511 https://llvm.org/docs/LangRef.html#extractelement-instruction 512 513 Example: 514 515 ```mlir 516 %c = arith.constant 15 : i32 517 %1 = vector.extractelement %0[%c : i32]: vector<16xf32> 518 %2 = vector.extractelement %z[]: vector<f32> 519 ``` 520 }]; 521 let assemblyFormat = [{ 522 $vector `[` ($position^ `:` type($position))? `]` attr-dict `:` type($vector) 523 }]; 524 525 let builders = [ 526 // 0-D builder. 527 OpBuilder<(ins "Value":$source)>, 528 // 1-D + position builder. 529 OpBuilder<(ins "Value":$source, "Value":$position)>, 530 ]; 531 let extraClassDeclaration = [{ 532 VectorType getVectorType() { 533 return getVector().getType().cast<VectorType>(); 534 } 535 }]; 536 let hasVerifier = 1; 537 let hasFolder = 1; 538} 539 540def Vector_ExtractOp : 541 Vector_Op<"extract", [NoSideEffect, 542 PredOpTrait<"operand and result have same element type", 543 TCresVTEtIsSameAsOpBase<0, 0>>, 544 DeclareOpInterfaceMethods<InferTypeOpInterface>]>, 545 Arguments<(ins AnyVector:$vector, I64ArrayAttr:$position)>, 546 Results<(outs AnyType)> { 547 let summary = "extract operation"; 548 let description = [{ 549 Takes an n-D vector and a k-D position and extracts the (n-k)-D vector at 550 the proper position. Degenerates to an element type in the 0-D case. 551 552 Example: 553 554 ```mlir 555 %1 = vector.extract %0[3]: vector<4x8x16xf32> 556 %2 = vector.extract %0[3, 3, 3]: vector<4x8x16xf32> 557 ``` 558 }]; 559 let builders = [ 560 OpBuilder<(ins "Value":$source, "ArrayRef<int64_t>":$position)>, 561 // Convenience builder which assumes the values in `position` are defined by 562 // ConstantIndexOp. 563 OpBuilder<(ins "Value":$source, "ValueRange":$position)> 564 ]; 565 let extraClassDeclaration = [{ 566 static StringRef getPositionAttrStrName() { return "position"; } 567 VectorType getVectorType() { 568 return getVector().getType().cast<VectorType>(); 569 } 570 static bool isCompatibleReturnTypes(TypeRange l, TypeRange r); 571 }]; 572 let assemblyFormat = "$vector `` $position attr-dict `:` type($vector)"; 573 let hasCanonicalizer = 1; 574 let hasFolder = 1; 575 let hasVerifier = 1; 576} 577 578def Vector_ExtractMapOp : 579 Vector_Op<"extract_map", [NoSideEffect]>, 580 Arguments<(ins AnyVector:$vector, Variadic<Index>:$ids)>, 581 Results<(outs AnyVector)> { 582 let summary = "vector extract map operation"; 583 let description = [{ 584 Takes an N-D vector and extracts a sub-part of the vector starting at id 585 along each dimension. 586 587 The dimension associated to each element of `ids` used to extract are 588 implicitly deduced from the destination type. For each dimension the 589 multiplicity is the destination dimension size divided by the source 590 dimension size, each dimension with a multiplicity greater than 1 is 591 associated to the next id, following ids order. 592 For example if the source type is `vector<64x4x32xf32>` and the destination 593 type is `vector<4x4x2xf32>`, the first id maps to dimension 0 and the second 594 id to dimension 2. 595 596 Similarly to vector.tuple_get, this operation is used for progressive 597 lowering and should be folded away before converting to LLVM. 598 599 It is different than `vector.extract_slice` and 600 `vector.extract_strided_slice` as it takes a Value as index instead of an 601 attribute. Also in the future it is meant to support extracting along any 602 dimensions and not only the most major ones. 603 604 For instance: 605 ``` 606 // dynamic computation producing the value 0 of index type 607 %idx0 = ... : index 608 // dynamic computation producing the value 1 of index type 609 %idx1 = ... : index 610 %0 = arith.constant dense<0, 1, 2, 3>: vector<4xi32> 611 // extracts values [0, 1] 612 %1 = vector.extract_map %0[%idx0] : vector<4xi32> to vector<2xi32> 613 // extracts values [1, 2] 614 %2 = vector.extract_map %0[%idx1] : vector<4xi32> to vector<2xi32> 615 ``` 616 617 Example: 618 619 ```mlir 620 %ev = vector.extract_map %v[%id] : vector<32xf32> to vector<1xf32> 621 %ev1 = vector.extract_map %v1[%id1, %id2] : vector<64x4x32xf32> 622 to vector<4x4x2xf32> 623 ``` 624 }]; 625 let builders = [ 626 OpBuilder<(ins "Value":$vector, "ValueRange":$ids, 627 "ArrayRef<int64_t>":$multiplicity, 628 "AffineMap":$map)>]; 629 let extraClassDeclaration = [{ 630 VectorType getSourceVectorType() { 631 return getVector().getType().cast<VectorType>(); 632 } 633 VectorType getResultType() { 634 return getResult().getType().cast<VectorType>(); 635 } 636 void getMultiplicity(SmallVectorImpl<int64_t> &multiplicity); 637 AffineMap map(); 638 }]; 639 let assemblyFormat = [{ 640 $vector `[` $ids `]` attr-dict `:` type($vector) `to` type(results) 641 }]; 642 643 let hasFolder = 1; 644 let hasVerifier = 1; 645} 646 647def Vector_FMAOp : 648 Op<Vector_Dialect, "fma", [ 649 NoSideEffect, AllTypesMatch<["lhs", "rhs", "acc", "result"]>, 650 DeclareOpInterfaceMethods<VectorUnrollOpInterface, ["getShapeForUnroll"]> 651 ] # ElementwiseMappable.traits>, 652 Arguments<(ins AnyVector:$lhs, AnyVector:$rhs, AnyVector:$acc)>, 653 Results<(outs AnyVector:$result)> { 654 let summary = "vector fused multiply-add"; 655 let description = [{ 656 Multiply-add expressions operate on n-D vectors and compute a fused 657 pointwise multiply-and-accumulate: `$result = `$lhs * $rhs + $acc`. 658 All operands and result have the same vector type. The semantics 659 of the operation correspond to those of the `llvm.fma` 660 [intrinsic](https://llvm.org/docs/LangRef.html#int-fma). In the 661 particular case of lowering to LLVM, this is guaranteed to lower 662 to the `llvm.fma.*` intrinsic. 663 664 Example: 665 666 ```mlir 667 %3 = vector.fma %0, %1, %2: vector<8x16xf32> 668 ``` 669 }]; 670 let assemblyFormat = "$lhs `,` $rhs `,` $acc attr-dict `:` type($lhs)"; 671 let extraClassDeclaration = [{ 672 VectorType getVectorType() { return getLhs().getType().cast<VectorType>(); } 673 }]; 674} 675 676def Vector_InsertElementOp : 677 Vector_Op<"insertelement", [NoSideEffect, 678 TypesMatchWith<"source operand type matches element type of result", 679 "result", "source", 680 "$_self.cast<ShapedType>().getElementType()">, 681 AllTypesMatch<["dest", "result"]>]>, 682 Arguments<(ins AnyType:$source, AnyVectorOfAnyRank:$dest, 683 Optional<AnySignlessIntegerOrIndex>:$position)>, 684 Results<(outs AnyVectorOfAnyRank:$result)> { 685 let summary = "insertelement operation"; 686 let description = [{ 687 Takes a scalar source, a 0-D or 1-D destination vector and a dynamic index 688 position and inserts the source into the destination at the proper position. 689 690 Note that this instruction resembles vector.insert, but is restricted to 0-D 691 and 1-D vectors and relaxed to dynamic indices. 692 693 It is meant to be closer to LLVM's version: 694 https://llvm.org/docs/LangRef.html#insertelement-instruction 695 696 Example: 697 698 ```mlir 699 %c = arith.constant 15 : i32 700 %f = arith.constant 0.0f : f32 701 %1 = vector.insertelement %f, %0[%c : i32]: vector<16xf32> 702 %2 = vector.insertelement %f, %z[]: vector<f32> 703 ``` 704 }]; 705 let assemblyFormat = [{ 706 $source `,` $dest `[` ($position^ `:` type($position))? `]` attr-dict `:` 707 type($result) 708 }]; 709 710 let builders = [ 711 // 0-D builder. 712 OpBuilder<(ins "Value":$source, "Value":$dest)>, 713 ]; 714 let extraClassDeclaration = [{ 715 Type getSourceType() { return getSource().getType(); } 716 VectorType getDestVectorType() { 717 return getDest().getType().cast<VectorType>(); 718 } 719 }]; 720 let hasVerifier = 1; 721 let hasFolder = 1; 722} 723 724def Vector_InsertOp : 725 Vector_Op<"insert", [NoSideEffect, 726 PredOpTrait<"source operand and result have same element type", 727 TCresVTEtIsSameAsOpBase<0, 0>>, 728 AllTypesMatch<["dest", "res"]>]>, 729 Arguments<(ins AnyType:$source, AnyVector:$dest, I64ArrayAttr:$position)>, 730 Results<(outs AnyVector:$res)> { 731 let summary = "insert operation"; 732 let description = [{ 733 Takes an n-D source vector, an (n+k)-D destination vector and a k-D position 734 and inserts the n-D source into the (n+k)-D destination at the proper 735 position. Degenerates to a scalar source type when n = 0. 736 737 Example: 738 739 ```mlir 740 %2 = vector.insert %0, %1[3] : vector<8x16xf32> into vector<4x8x16xf32> 741 %5 = vector.insert %3, %4[3, 3, 3] : f32 into vector<4x8x16xf32> 742 ``` 743 }]; 744 let assemblyFormat = [{ 745 $source `,` $dest $position attr-dict `:` type($source) `into` type($dest) 746 }]; 747 748 let builders = [ 749 OpBuilder<(ins "Value":$source, "Value":$dest, 750 "ArrayRef<int64_t>":$position)>, 751 // Convenience builder which assumes all values are constant indices. 752 OpBuilder<(ins "Value":$source, "Value":$dest, "ValueRange":$position)> 753 ]; 754 let extraClassDeclaration = [{ 755 static StringRef getPositionAttrStrName() { return "position"; } 756 Type getSourceType() { return getSource().getType(); } 757 VectorType getDestVectorType() { 758 return getDest().getType().cast<VectorType>(); 759 } 760 }]; 761 762 let hasCanonicalizer = 1; 763 let hasFolder = 1; 764 let hasVerifier = 1; 765} 766 767def Vector_InsertMapOp : 768 Vector_Op<"insert_map", [NoSideEffect, AllTypesMatch<["dest", "result"]>]>, 769 Arguments<(ins AnyVector:$vector, AnyVector:$dest, Variadic<Index>:$ids)>, 770 Results<(outs AnyVector:$result)> { 771 let summary = "vector insert map operation"; 772 let description = [{ 773 Inserts a N-D vector and within a larger vector starting at id. The new 774 vector created will have the same size as the destination operand vector. 775 776 The dimension associated to each element of `ids` used to insert is 777 implicitly deduced from the source type (see `ExtractMapOp` for details). 778 For example if source type is `vector<4x4x2xf32>` and the destination type 779 is `vector<64x4x32xf32>`, the first id maps to dimension 0 and the second id 780 to dimension 2. 781 782 Similarly to vector.tuple_get, this operation is used for progressive 783 lowering and should be folded away before converting to LLVM. 784 785 It is different than `vector.insert` and `vector.insert_strided_slice` as it 786 takes a Value as index instead of an attribute. Also in the future it is 787 meant to support inserting along any dimensions and not only the most major 788 ones. 789 790 This operations is meant to be used in combination with vector.extract_map. 791 792 For instance: 793 ``` 794 // dynamic computation producing the value 0 of index type 795 %idx0 = ... : index 796 // dynamic computation producing the value 1 of index type 797 %idx1 = ... : index / 798 %0 = arith.constant dense<0, 1, 2, 3>: vector<4xi32> 799 // extracts values [0, 1] 800 %1 = vector.extract_map %0[%idx0] : vector<4xi32> to vector<2xi32> 801 // extracts values [1, 2] 802 %2 = vector.extract_map %0[%idx1] : vector<4xi32> to vector<2xi32> 803 // insert [0, 1] into [x, x, x, x] and produce [0, 1, x, x] 804 %3 = vector.insert_map %1, %0[%idx0] : vector<2xi32> into vector<4xi32> 805 // insert [1, 2] into [x, x, x, x] and produce [x, 1, 2, x] 806 %4 = vector.insert_map %2, %0[%idx1] : vector<2xi32> into vector<4xi32> 807 ``` 808 Example: 809 810 ```mlir 811 %v = vector.insert_map %ev %v[%id] : vector<1xf32> into vector<32xf32> 812 %v1 = vector.insert_map %ev1, %v1[%arg0, %arg1] : vector<2x4x1xf32> 813 into vector<64x4x32xf32> 814 ``` 815 }]; 816 let extraClassDeclaration = [{ 817 VectorType getSourceVectorType() { 818 return getVector().getType().cast<VectorType>(); 819 } 820 VectorType getResultType() { 821 return getResult().getType().cast<VectorType>(); 822 } 823 // Return a map indicating the dimension mapping to the given ids. 824 AffineMap map(); 825 }]; 826 let assemblyFormat = [{ 827 $vector `,` $dest `[` $ids `]` attr-dict 828 `:` type($vector) `into` type($result) 829 }]; 830 let hasVerifier = 1; 831} 832 833def Vector_InsertStridedSliceOp : 834 Vector_Op<"insert_strided_slice", [NoSideEffect, 835 PredOpTrait<"operand #0 and result have same element type", 836 TCresVTEtIsSameAsOpBase<0, 0>>, 837 AllTypesMatch<["dest", "res"]>]>, 838 Arguments<(ins AnyVector:$source, AnyVector:$dest, I64ArrayAttr:$offsets, 839 I64ArrayAttr:$strides)>, 840 Results<(outs AnyVector:$res)> { 841 let summary = "strided_slice operation"; 842 let description = [{ 843 Takes a k-D source vector, an n-D destination vector (n >= k), n-sized 844 `offsets` integer array attribute, a k-sized `strides` integer array attribute 845 and inserts the k-D source vector as a strided subvector at the proper offset 846 into the n-D destination vector. 847 848 At the moment strides must contain only 1s. 849 850 Returns an n-D vector that is a copy of the n-D destination vector in which 851 the last k-D dimensions contain the k-D source vector elements strided at 852 the proper location as specified by the offsets. 853 854 Example: 855 856 ```mlir 857 %2 = vector.insert_strided_slice %0, %1 858 {offsets = [0, 0, 2], strides = [1, 1]}: 859 vector<2x4xf32> into vector<16x4x8xf32> 860 ``` 861 }]; 862 863 let assemblyFormat = [{ 864 $source `,` $dest attr-dict `:` type($source) `into` type($dest) 865 }]; 866 867 let builders = [ 868 OpBuilder<(ins "Value":$source, "Value":$dest, 869 "ArrayRef<int64_t>":$offsets, "ArrayRef<int64_t>":$strides)> 870 ]; 871 let extraClassDeclaration = [{ 872 static StringRef getOffsetsAttrStrName() { return "offsets"; } 873 static StringRef getStridesAttrStrName() { return "strides"; } 874 VectorType getSourceVectorType() { 875 return getSource().getType().cast<VectorType>(); 876 } 877 VectorType getDestVectorType() { 878 return getDest().getType().cast<VectorType>(); 879 } 880 bool hasNonUnitStrides() { 881 return llvm::any_of(getStrides(), [](Attribute attr) { 882 return attr.cast<IntegerAttr>().getInt() != 1; 883 }); 884 } 885 }]; 886 887 let hasFolder = 1; 888 let hasVerifier = 1; 889 let hasCanonicalizer = 1; 890} 891 892def Vector_OuterProductOp : 893 Vector_Op<"outerproduct", [NoSideEffect, 894 PredOpTrait<"lhs operand and result have same element type", 895 TCresVTEtIsSameAsOpBase<0, 0>>, 896 PredOpTrait<"rhs operand and result have same element type", 897 TCresVTEtIsSameAsOpBase<0, 1>>]>, 898 Arguments<(ins AnyVector:$lhs, AnyType:$rhs, 899 Variadic<AnyVector>:$acc, 900 DefaultValuedAttr<Vector_CombiningKindAttr, "CombiningKind::ADD">:$kind)>, 901 Results<(outs AnyVector)> { 902 let summary = "vector outerproduct with optional fused add"; 903 let description = [{ 904 Takes 2 1-D vectors and returns the 2-D vector containing the outer-product, 905 as illustrated below: 906 ``` 907 outer | [c, d] 908 ------+------------ 909 [a, | [ [a*c, a*d], 910 b] | [b*c, b*d] ] 911 ``` 912 This operation also accepts a 1-D vector lhs and a scalar rhs. In this 913 case a simple AXPY operation is performed, which returns a 1-D vector. 914 ``` 915 [a, b] * c = [a*c, b*c] 916 ``` 917 918 An optional extra vector argument with the same shape as the output 919 vector may be specified in which case the operation returns the sum of 920 the outer-product and the extra vector. In this multiply-accumulate 921 scenario for floating-point arguments, the rounding mode is enforced 922 by guaranteeing that a fused-multiply add operation is emitted. When 923 lowered to the LLVMIR dialect, this form emits `llvm.intr.fma`, which 924 is guaranteed to lower to actual `fma` instructions on x86. 925 926 An optional kind attribute may be specified to be add/mul/min/max 927 for int/fp, and and/or/xor for int only. The default is "add", in which 928 case the operation returns a fused multiply-add. In other cases it returns 929 a multiply followed by the appropriate operation (for example, a compare and 930 select for "max"). 931 932 Example: 933 934 ``` 935 %2 = vector.outerproduct %0, %1: vector<4xf32>, vector<8xf32> 936 return %2: vector<4x8xf32> 937 938 %3 = vector.outerproduct %0, %1, %2: 939 vector<4xf32>, vector<8xf32>, vector<4x8xf32> 940 return %3: vector<4x8xf32> 941 942 %4 = vector.outerproduct %0, %1, %2 {kind = #vector.kind<max>}: 943 vector<4xf32>, vector<8xf32>, vector<4x8xf32> 944 return %3: vector<4x8xf32> 945 946 %6 = vector.outerproduct %4, %5: vector<10xf32>, f32 947 return %6: vector<10xf32> 948 949 ``` 950 }]; 951 let builders = [ 952 // Build an op without mask, use the type of `acc` as the return type. 953 OpBuilder<(ins "Value":$lhs, "Value":$rhs, "Value":$acc)> 954 ]; 955 let extraClassDeclaration = [{ 956 VectorType getOperandVectorTypeLHS() { 957 return getLhs().getType().cast<VectorType>(); 958 } 959 Type getOperandTypeRHS() { 960 return getRhs().getType(); 961 } 962 VectorType getOperandVectorTypeACC() { 963 return (llvm::size(getAcc()) == 0) 964 ? VectorType() 965 : (*getAcc().begin()).getType().cast<VectorType>(); 966 } 967 VectorType getVectorType() { 968 return getResult().getType().cast<VectorType>(); 969 } 970 static constexpr StringRef getKindAttrStrName() { 971 return "kind"; 972 } 973 static CombiningKind getDefaultKind() { 974 return CombiningKind::ADD; 975 } 976 }]; 977 let hasCustomAssemblyFormat = 1; 978 let hasVerifier = 1; 979} 980 981// TODO: Add transformation which decomposes ReshapeOp into an optimized 982// sequence of vector rotate/shuffle/select operations. 983def Vector_ReshapeOp : 984 Vector_Op<"reshape", [AttrSizedOperandSegments, NoSideEffect]>, 985 Arguments<(ins AnyVector:$vector, Variadic<Index>:$input_shape, 986 Variadic<Index>:$output_shape, 987 I64ArrayAttr:$fixed_vector_sizes)>, 988 Results<(outs AnyVector:$result)> { 989 let summary = "vector reshape operation"; 990 let description = [{ 991 Reshapes its vector operand from 'input_shape' to 'output_shape' maintaining 992 fixed vector dimension 'fixed_vector_sizes' on the innermost vector 993 dimensions. 994 995 The parameters 'input_shape' and 'output_shape' represent valid data shapes 996 across fixed vector shapes. For example, if a vector has a valid data 997 shape [6] with fixed vector size [8], then the valid data elements are 998 assumed to be stored at the beginning of the vector with the remaining 999 vector elements undefined. 1000 1001 In the examples below, valid data elements are represented by an alphabetic 1002 character, and undefined data elements are represented by '-'. 1003 1004 Example 1005 1006 vector<1x8xf32> with valid data shape [6], fixed vector sizes [8] 1007 1008 input: [a, b, c, d, e, f] 1009 1010 layout map: (d0) -> (d0 floordiv 8, d0 mod 8) 1011 1012 vector layout: [a, b, c, d, e, f, -, -] 1013 1014 Example 1015 1016 vector<2x8xf32> with valid data shape [10], fixed vector sizes [8] 1017 1018 input: [a, b, c, d, e, f, g, h, i, j] 1019 1020 layout map: (d0) -> (d0 floordiv 8, d0 mod 8) 1021 1022 vector layout: [[a, b, c, d, e, f, g, h], 1023 [i, j, -, -, -, -, -, -]] 1024 1025 Example 1026 1027 vector<2x2x2x3xf32> with valid data shape [3, 5], fixed vector sizes 1028 [2, 3] 1029 1030 input: [[a, b, c, d, e], 1031 [f, g, h, i, j], 1032 [k, l, m, n, o]] 1033 1034 layout map: (d0, d1) -> (d0 floordiv 3, d1 floordiv 5, 1035 d0 mod 3, d1 mod 5) 1036 1037 vector layout: [[[[a, b, c], 1038 [f, g, h]] 1039 [[d, e, -], 1040 [i, j, -]]], 1041 [[[k, l, m], 1042 [-, -, -]] 1043 [[n, o, -], 1044 [-, -, -]]]] 1045 1046 Example 1047 1048 %1 = vector.reshape %0, [%c3, %c6], [%c2, %c9], [4] 1049 : vector<3x2x4xf32> to vector<2x3x4xf32> 1050 1051 input: [[a, b, c, d, e, f], 1052 [g, h, i, j, k, l], 1053 [m, n, o, p, q, r]] 1054 1055 layout map: (d0, d1) -> (d0, d1 floordiv 4, d1 mod 4) 1056 1057 1058 Input vector: [[[a, b, c, d], 1059 [e, f, -, -]], 1060 [[g, h, i, j], 1061 [k, l, -, -]], 1062 [[m, n, o, p], 1063 [q, r, -, -]]] 1064 1065 Output vector: [[[a, b, c, d], 1066 [e, f, g, h], 1067 [i, -, -, -]], 1068 [[j, k, l, m], 1069 [n, o, p, q], 1070 [r, -, -, -]]] 1071 }]; 1072 1073 let extraClassDeclaration = [{ 1074 VectorType getInputVectorType() { 1075 return getVector().getType().cast<VectorType>(); 1076 } 1077 VectorType getOutputVectorType() { 1078 return getResult().getType().cast<VectorType>(); 1079 } 1080 1081 /// Returns as integer value the number of input shape operands. 1082 int64_t getNumInputShapeSizes() { return getInputShape().size(); } 1083 1084 /// Returns as integer value the number of output shape operands. 1085 int64_t getNumOutputShapeSizes() { return getOutputShape().size(); } 1086 1087 void getFixedVectorSizes(SmallVectorImpl<int64_t> &results); 1088 1089 static StringRef getFixedVectorSizesAttrStrName() { 1090 return "fixed_vector_sizes"; 1091 } 1092 static StringRef getInputShapeAttrStrName() { return "input_shape"; } 1093 static StringRef getOutputShapeAttrStrName() { return "output_shape"; } 1094 }]; 1095 1096 let assemblyFormat = [{ 1097 $vector `,` `[` $input_shape `]` `,` `[` $output_shape `]` `,` 1098 $fixed_vector_sizes attr-dict `:` type($vector) `to` type($result) 1099 }]; 1100 let hasVerifier = 1; 1101} 1102 1103def Vector_ExtractStridedSliceOp : 1104 Vector_Op<"extract_strided_slice", [NoSideEffect, 1105 PredOpTrait<"operand and result have same element type", 1106 TCresVTEtIsSameAsOpBase<0, 0>>]>, 1107 Arguments<(ins AnyVector:$vector, I64ArrayAttr:$offsets, 1108 I64ArrayAttr:$sizes, I64ArrayAttr:$strides)>, 1109 Results<(outs AnyVector)> { 1110 let summary = "extract_strided_slice operation"; 1111 let description = [{ 1112 Takes an n-D vector, k-D `offsets` integer array attribute, a k-sized 1113 `sizes` integer array attribute, a k-sized `strides` integer array 1114 attribute and extracts the n-D subvector at the proper offset. 1115 1116 At the moment strides must contain only 1s. 1117 // TODO: support non-1 strides. 1118 1119 Returns an n-D vector where the first k-D dimensions match the `sizes` 1120 attribute. The returned subvector contains the elements starting at offset 1121 `offsets` and ending at `offsets + sizes`. 1122 1123 Example: 1124 1125 ```mlir 1126 %1 = vector.extract_strided_slice %0 1127 {offsets = [0, 2], sizes = [2, 4], strides = [1, 1]}: 1128 vector<4x8x16xf32> to vector<2x4x16xf32> 1129 1130 // TODO: Evolve to a range form syntax similar to: 1131 %1 = vector.extract_strided_slice %0[0:2:1][2:4:1] 1132 vector<4x8x16xf32> to vector<2x4x16xf32> 1133 ``` 1134 }]; 1135 let builders = [ 1136 OpBuilder<(ins "Value":$source, "ArrayRef<int64_t>":$offsets, 1137 "ArrayRef<int64_t>":$sizes, "ArrayRef<int64_t>":$strides)> 1138 ]; 1139 let extraClassDeclaration = [{ 1140 static StringRef getOffsetsAttrStrName() { return "offsets"; } 1141 static StringRef getSizesAttrStrName() { return "sizes"; } 1142 static StringRef getStridesAttrStrName() { return "strides"; } 1143 VectorType getVectorType(){ return getVector().getType().cast<VectorType>(); } 1144 void getOffsets(SmallVectorImpl<int64_t> &results); 1145 bool hasNonUnitStrides() { 1146 return llvm::any_of(getStrides(), [](Attribute attr) { 1147 return attr.cast<IntegerAttr>().getInt() != 1; 1148 }); 1149 } 1150 }]; 1151 let hasCanonicalizer = 1; 1152 let hasFolder = 1; 1153 let hasVerifier = 1; 1154 let assemblyFormat = "$vector attr-dict `:` type($vector) `to` type(results)"; 1155} 1156 1157def Vector_TransferReadOp : 1158 Vector_Op<"transfer_read", [ 1159 DeclareOpInterfaceMethods<VectorTransferOpInterface>, 1160 DeclareOpInterfaceMethods<VectorUnrollOpInterface, ["getShapeForUnroll"]>, 1161 DeclareOpInterfaceMethods<MemoryEffectsOpInterface>, 1162 AttrSizedOperandSegments 1163 ]>, 1164 Arguments<(ins AnyShaped:$source, 1165 Variadic<Index>:$indices, 1166 AffineMapAttr:$permutation_map, 1167 AnyType:$padding, 1168 Optional<VectorOf<[I1]>>:$mask, 1169 OptionalAttr<BoolArrayAttr>:$in_bounds)>, 1170 Results<(outs AnyVectorOfAnyRank:$vector)> { 1171 1172 let summary = "Reads a supervector from memory into an SSA vector value."; 1173 1174 let description = [{ 1175 The `vector.transfer_read` op performs a read from a slice within a 1176 [MemRef](../LangRef.md#memref-type) or a Ranked 1177 [Tensor](../LangRef.md#tensor-type) supplied as its first operand 1178 into a [vector](../LangRef.md#vector-type) of the same base elemental type. 1179 1180 A memref/tensor operand with vector element type, must have its vector 1181 element type match a suffix (shape and element type) of the vector (e.g. 1182 memref<3x2x6x4x3xf32>, vector<1x1x4x3xf32>). 1183 1184 The slice is further defined by a full-rank index within the MemRef/Tensor, 1185 supplied as the operands `[1 .. 1 + rank(memref/tensor))` that defines the 1186 starting point of the transfer (e.g. `%A[%i0, %i1, %i2]`). 1187 1188 The permutation_map [attribute](../LangRef.md#attributes) is an 1189 [affine-map](Affine.md#affine-maps) which specifies the transposition on the 1190 slice to match the vector shape. The permutation map may be implicit and 1191 omitted from parsing and printing if it is the canonical minor identity map 1192 (i.e. if it does not permute or broadcast any dimension). 1193 1194 The size of the slice is specified by the size of the vector, given as the 1195 return type. 1196 1197 An SSA value `padding` of the same elemental type as the MemRef/Tensor is 1198 provided to specify a fallback value in the case of out-of-bounds accesses 1199 and/or masking. 1200 1201 An optional SSA value `mask` of the same shape as the vector type may be 1202 specified to mask out elements. Such elements will be replaces with 1203 `padding`. Elements whose corresponding mask element is `0` are masked out. 1204 1205 An optional boolean array attribute `in_bounds` specifies for every vector 1206 dimension if the transfer is guaranteed to be within the source bounds. 1207 While the starting point of the transfer has to be in-bounds, accesses may 1208 run out-of-bounds as indices increase. Broadcast dimensions must always be 1209 in-bounds. If specified, the `in_bounds` array length has to be equal to the 1210 vector rank. In absence of the attribute, accesses along all dimensions 1211 (except for broadcasts) may run out-of-bounds. A `vector.transfer_read` can 1212 be lowered to a simple load if all dimensions are specified to be within 1213 bounds and no `mask` was specified. 1214 1215 This operation is called 'read' by opposition to 'load' because the 1216 super-vector granularity is generally not representable with a single 1217 hardware register. A `vector.transfer_read` is thus a mid-level abstraction 1218 that supports super-vectorization with non-effecting padding for full-tile 1219 only operations. 1220 1221 More precisely, let's dive deeper into the permutation_map for the following 1222 MLIR: 1223 1224 ```mlir 1225 vector.transfer_read %A[%expr1, %expr2, %expr3, %expr4] 1226 { permutation_map : (d0,d1,d2,d3) -> (d2,0,d0) } : 1227 memref<?x?x?x?xf32>, vector<3x4x5xf32> 1228 ``` 1229 1230 This operation always reads a slice starting at `%A[%expr1, %expr2, %expr3, 1231 %expr4]`. The size of the slice is 3 along d2 and 5 along d0, so the slice 1232 is: `%A[%expr1 : %expr1 + 5, %expr2, %expr3:%expr3 + 3, %expr4]` 1233 1234 That slice needs to be read into a `vector<3x4x5xf32>`. Since the 1235 permutation map is not full rank, there must be a broadcast along vector 1236 dimension `1`. 1237 1238 A notional lowering of vector.transfer_read could generate code resembling: 1239 1240 ```mlir 1241 // %expr1, %expr2, %expr3, %expr4 defined before this point 1242 %tmp = alloc() : vector<3x4x5xf32> 1243 %view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<3x4x5xf32>> 1244 for %i = 0 to 3 { 1245 affine.for %j = 0 to 4 { 1246 affine.for %k = 0 to 5 { 1247 %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] : 1248 memref<?x?x?x?xf32> 1249 store %tmp[%i, %j, %k] : vector<3x4x5xf32> 1250 }}} 1251 %c0 = arith.constant 0 : index 1252 %vec = load %view_in_tmp[%c0] : vector<3x4x5xf32> 1253 ``` 1254 1255 On a GPU one could then map `i`, `j`, `k` to blocks and threads. Notice that 1256 the temporary storage footprint is `3 * 5` values but `3 * 4 * 5` values are 1257 actually transferred between `%A` and `%tmp`. 1258 1259 Alternatively, if a notional vector broadcast operation were available, the 1260 lowered code would resemble: 1261 1262 ```mlir 1263 // %expr1, %expr2, %expr3, %expr4 defined before this point 1264 %tmp = alloc() : vector<3x4x5xf32> 1265 %view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<3x4x5xf32>> 1266 for %i = 0 to 3 { 1267 affine.for %k = 0 to 5 { 1268 %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] : 1269 memref<?x?x?x?xf32> 1270 store %tmp[%i, 0, %k] : vector<3x4x5xf32> 1271 }} 1272 %c0 = arith.constant 0 : index 1273 %tmpvec = load %view_in_tmp[%c0] : vector<3x4x5xf32> 1274 %vec = broadcast %tmpvec, 1 : vector<3x4x5xf32> 1275 ``` 1276 1277 where `broadcast` broadcasts from element 0 to all others along the 1278 specified dimension. This time, the temporary storage footprint is `3 * 5` 1279 values which is the same amount of data as the `3 * 5` values transferred. 1280 An additional `1` broadcast is required. On a GPU this broadcast could be 1281 implemented using a warp-shuffle if loop `j` were mapped to `threadIdx.x`. 1282 1283 Syntax 1284 ``` 1285 operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list 1286 `{` attribute-entry `} :` memref-type `,` vector-type 1287 ``` 1288 1289 Example: 1290 1291 ```mlir 1292 // Read the slice `%A[%i0, %i1:%i1+256, %i2:%i2+32]` into vector<32x256xf32> 1293 // and pad with %f0 to handle the boundary case: 1294 %f0 = arith.constant 0.0f : f32 1295 for %i0 = 0 to %0 { 1296 affine.for %i1 = 0 to %1 step 256 { 1297 affine.for %i2 = 0 to %2 step 32 { 1298 %v = vector.transfer_read %A[%i0, %i1, %i2], (%f0) 1299 {permutation_map: (d0, d1, d2) -> (d2, d1)} : 1300 memref<?x?x?xf32>, vector<32x256xf32> 1301 }}} 1302 1303 // Read the slice `%A[%i0, %i1]` (i.e. the element `%A[%i0, %i1]`) into 1304 // vector<128xf32>. The underlying implementation will require a 1-D vector 1305 // broadcast: 1306 for %i0 = 0 to %0 { 1307 affine.for %i1 = 0 to %1 { 1308 %3 = vector.transfer_read %A[%i0, %i1] 1309 {permutation_map: (d0, d1) -> (0)} : 1310 memref<?x?xf32>, vector<128xf32> 1311 } 1312 } 1313 1314 // Read from a memref with vector element type. 1315 %4 = vector.transfer_read %arg1[%c3, %c3], %vf0 1316 {permutation_map = (d0, d1)->(d0, d1)} 1317 : memref<?x?xvector<4x3xf32>>, vector<1x1x4x3xf32> 1318 1319 // Read from a tensor with vector element type. 1320 %4 = vector.transfer_read %arg1[%c3, %c3], %vf0 1321 {permutation_map = (d0, d1)->(d0, d1)} 1322 : tensor<?x?xvector<4x3xf32>>, vector<1x1x4x3xf32> 1323 1324 // Special encoding for 0-d transfer with 0-d tensor/memref, vector shape 1325 // {1} and permutation_map () -> (0). 1326 %0 = vector.transfer_read %arg0[], %f0 {permutation_map = affine_map<()->(0)>} : 1327 tensor<f32>, vector<1xf32> 1328 ``` 1329 }]; 1330 1331 let builders = [ 1332 /// 1. Builder that sets padding to zero and an empty mask (variant with attrs). 1333 OpBuilder<(ins "VectorType":$vectorType, 1334 "Value":$source, 1335 "ValueRange":$indices, 1336 "AffineMapAttr":$permutationMapAttr, 1337 "ArrayAttr":$inBoundsAttr)>, 1338 /// 2. Builder that sets padding to zero and an empty mask (variant without attrs). 1339 OpBuilder<(ins "VectorType":$vectorType, 1340 "Value":$source, 1341 "ValueRange":$indices, 1342 "AffineMap":$permutationMap, 1343 CArg<"Optional<ArrayRef<bool>>", "::llvm::None">:$inBounds)>, 1344 /// 3. Builder that sets permutation map to 'getMinorIdentityMap'. 1345 OpBuilder<(ins "VectorType":$vectorType, 1346 "Value":$source, 1347 "ValueRange":$indices, 1348 "Value":$padding, 1349 CArg<"Optional<ArrayRef<bool>>", "::llvm::None">:$inBounds)>, 1350 /// 4. Builder that sets padding to zero and permutation map to 1351 /// 'getMinorIdentityMap'. 1352 OpBuilder<(ins "VectorType":$vectorType, 1353 "Value":$source, 1354 "ValueRange":$indices, 1355 CArg<"Optional<ArrayRef<bool>>", "::llvm::None">:$inBounds)>, 1356 ]; 1357 let hasCanonicalizer = 1; 1358 let hasCustomAssemblyFormat = 1; 1359 let hasFolder = 1; 1360 let hasVerifier = 1; 1361} 1362 1363def Vector_TransferWriteOp : 1364 Vector_Op<"transfer_write", [ 1365 DeclareOpInterfaceMethods<VectorTransferOpInterface>, 1366 DeclareOpInterfaceMethods<VectorUnrollOpInterface, ["getShapeForUnroll"]>, 1367 DeclareOpInterfaceMethods<MemoryEffectsOpInterface>, 1368 AttrSizedOperandSegments 1369 ]>, 1370 Arguments<(ins AnyVectorOfAnyRank:$vector, 1371 AnyShaped:$source, 1372 Variadic<Index>:$indices, 1373 AffineMapAttr:$permutation_map, 1374 Optional<VectorOf<[I1]>>:$mask, 1375 OptionalAttr<BoolArrayAttr>:$in_bounds)>, 1376 Results<(outs Optional<AnyRankedTensor>:$result)> { 1377 1378 let summary = "The vector.transfer_write op writes a supervector to memory."; 1379 1380 let description = [{ 1381 The `vector.transfer_write` op performs a write from a 1382 [vector](../LangRef.md#vector-type), supplied as its first operand, into a 1383 slice within a [MemRef](../LangRef.md#memref-type) or a Ranked 1384 [Tensor](../LangRef.md#tensor-type) of the same base elemental type, 1385 supplied as its second operand. 1386 1387 A vector memref/tensor operand must have its vector element type match a 1388 suffix (shape and element type) of the vector (e.g. memref<3x2x6x4x3xf32>, 1389 vector<1x1x4x3xf32>). If the operand is a tensor, the operation returns a 1390 new tensor of the same type. 1391 1392 The slice is further defined by a full-rank index within the MemRef/Tensor, 1393 supplied as the operands `[2 .. 2 + rank(memref/tensor))` that defines the 1394 starting point of the transfer (e.g. `%A[%i0, %i1, %i2, %i3]`). 1395 1396 The permutation_map [attribute](../LangRef.md#attributes) is an 1397 [affine-map](Affine.md#affine-maps) which specifies the transposition on the 1398 slice to match the vector shape. The permutation map may be implicit and 1399 omitted from parsing and printing if it is the canonical minor identity map 1400 (i.e. if it does not permute any dimension). In contrast to `transfer_read`, 1401 write ops cannot have broadcast dimensions. 1402 1403 The size of the slice is specified by the size of the vector. 1404 1405 An optional SSA value `mask` of the same shape as the vector type may be 1406 specified to mask out elements. Elements whose corresponding mask element 1407 is `0` are masked out. 1408 1409 An optional boolean array attribute `in_bounds` specifies for every vector 1410 dimension if the transfer is guaranteed to be within the source bounds. 1411 While the starting point of the transfer has to be in-bounds, accesses may 1412 run out-of-bounds as indices increase. If specified, the `in_bounds` array 1413 length has to be equal to the vector rank. In absence of the attribute, 1414 accesses along all dimensions may run out-of-bounds. A 1415 `vector.transfer_write` can be lowered to a simple store if all dimensions 1416 are specified to be within bounds and no `mask` was specified. 1417 1418 This operation is called 'write' by opposition to 'store' because the 1419 super-vector granularity is generally not representable with a single 1420 hardware register. A `vector.transfer_write` is thus a 1421 mid-level abstraction that supports super-vectorization with non-effecting 1422 padding for full-tile-only code. It is the responsibility of 1423 `vector.transfer_write`'s implementation to ensure the memory writes are 1424 valid. Different lowerings may be pertinent depending on the hardware 1425 support. 1426 1427 Example: 1428 1429 ```mlir 1430 // write vector<16x32x64xf32> into the slice 1431 // `%A[%i0, %i1:%i1+32, %i2:%i2+64, %i3:%i3+16]`: 1432 for %i0 = 0 to %0 { 1433 affine.for %i1 = 0 to %1 step 32 { 1434 affine.for %i2 = 0 to %2 step 64 { 1435 affine.for %i3 = 0 to %3 step 16 { 1436 %val = `ssa-value` : vector<16x32x64xf32> 1437 vector.transfer_write %val, %A[%i0, %i1, %i2, %i3] 1438 {permutation_map: (d0, d1, d2, d3) -> (d3, d1, d2)} : 1439 vector<16x32x64xf32>, memref<?x?x?x?xf32> 1440 }}}} 1441 1442 // write to a memref with vector element type. 1443 vector.transfer_write %4, %arg1[%c3, %c3] 1444 {permutation_map = (d0, d1)->(d0, d1)} 1445 : vector<1x1x4x3xf32>, memref<?x?xvector<4x3xf32>> 1446 1447 // return a tensor where the vector is inserted into the source tensor. 1448 %5 = vector.transfer_write %4, %arg1[%c3, %c3] 1449 {permutation_map = (d0, d1)->(d0, d1)} 1450 : vector<1x1x4x3xf32>, tensor<?x?xvector<4x3xf32>> 1451 1452 // Special encoding for 0-d transfer with 0-d tensor/memref, vector shape 1453 // {1} and permutation_map () -> (0). 1454 %1 = vector.transfer_write %0, %arg0[] {permutation_map = affine_map<()->(0)>} : 1455 vector<1xf32>, tensor<f32> 1456 ``` 1457 }]; 1458 1459 let builders = [ 1460 /// 1. Builder with type inference. 1461 OpBuilder<(ins "Value":$vector, 1462 "Value":$dest, 1463 "ValueRange":$indices, 1464 "AffineMapAttr":$permutationMapAttr, 1465 "Value":$mask, 1466 "ArrayAttr":$inBoundsAttr)>, 1467 /// 2. Builder with type inference that sets an empty mask (variant with attrs). 1468 OpBuilder<(ins "Value":$vector, 1469 "Value":$dest, 1470 "ValueRange":$indices, 1471 "AffineMapAttr":$permutationMapAttr, 1472 "ArrayAttr":$inBoundsAttr)>, 1473 /// 3. Builder with type inference that sets an empty mask (variant without attrs). 1474 OpBuilder<(ins "Value":$vector, 1475 "Value":$dest, 1476 "ValueRange":$indices, 1477 "AffineMap":$permutationMap, 1478 CArg<"Optional<ArrayRef<bool>>", "::llvm::None">:$inBounds)>, 1479 /// 4. Builder with type inference that sets an empty mask and sets permutation 1480 /// map to 'getMinorIdentityMap'. 1481 OpBuilder<(ins "Value":$vector, 1482 "Value":$dest, 1483 "ValueRange":$indices, 1484 CArg<"Optional<ArrayRef<bool>>", "::llvm::None">:$inBounds)>, 1485 ]; 1486 let hasFolder = 1; 1487 let hasCanonicalizer = 1; 1488 let hasCustomAssemblyFormat = 1; 1489 let hasVerifier = 1; 1490} 1491 1492def Vector_LoadOp : Vector_Op<"load"> { 1493 let summary = "reads an n-D slice of memory into an n-D vector"; 1494 let description = [{ 1495 The 'vector.load' operation reads an n-D slice of memory into an n-D 1496 vector. It takes a 'base' memref, an index for each memref dimension and a 1497 result vector type as arguments. It returns a value of the result vector 1498 type. The 'base' memref and indices determine the start memory address from 1499 which to read. Each index provides an offset for each memref dimension 1500 based on the element type of the memref. The shape of the result vector 1501 type determines the shape of the slice read from the start memory address. 1502 The elements along each dimension of the slice are strided by the memref 1503 strides. Only unit strides are allowed along the most minor memref 1504 dimension. These constraints guarantee that elements read along the first 1505 dimension of the slice are contiguous in memory. 1506 1507 The memref element type can be a scalar or a vector type. If the memref 1508 element type is a scalar, it should match the element type of the result 1509 vector. If the memref element type is vector, it should match the result 1510 vector type. 1511 1512 Example 1: 1-D vector load on a scalar memref. 1513 ```mlir 1514 %result = vector.load %base[%i, %j] : memref<100x100xf32>, vector<8xf32> 1515 ``` 1516 1517 Example 2: 1-D vector load on a vector memref. 1518 ```mlir 1519 %result = vector.load %memref[%i, %j] : memref<200x100xvector<8xf32>>, vector<8xf32> 1520 ``` 1521 1522 Example 3: 2-D vector load on a scalar memref. 1523 ```mlir 1524 %result = vector.load %memref[%i, %j] : memref<200x100xf32>, vector<4x8xf32> 1525 ``` 1526 1527 Example 4: 2-D vector load on a vector memref. 1528 ```mlir 1529 %result = vector.load %memref[%i, %j] : memref<200x100xvector<4x8xf32>>, vector<4x8xf32> 1530 ``` 1531 1532 Representation-wise, the 'vector.load' operation permits out-of-bounds 1533 reads. Support and implementation of out-of-bounds vector loads is 1534 target-specific. No assumptions should be made on the value of elements 1535 loaded out of bounds. Not all targets may support out-of-bounds vector 1536 loads. 1537 1538 Example 5: Potential out-of-bound vector load. 1539 ```mlir 1540 %result = vector.load %memref[%index] : memref<?xf32>, vector<8xf32> 1541 ``` 1542 1543 Example 6: Explicit out-of-bound vector load. 1544 ```mlir 1545 %result = vector.load %memref[%c0] : memref<7xf32>, vector<8xf32> 1546 ``` 1547 }]; 1548 1549 let arguments = (ins Arg<AnyMemRef, "the reference to load from", 1550 [MemRead]>:$base, 1551 Variadic<Index>:$indices); 1552 let results = (outs AnyVector:$result); 1553 1554 let extraClassDeclaration = [{ 1555 MemRefType getMemRefType() { 1556 return getBase().getType().cast<MemRefType>(); 1557 } 1558 1559 VectorType getVectorType() { 1560 return getResult().getType().cast<VectorType>(); 1561 } 1562 }]; 1563 1564 let hasFolder = 1; 1565 let hasVerifier = 1; 1566 1567 let assemblyFormat = 1568 "$base `[` $indices `]` attr-dict `:` type($base) `,` type($result)"; 1569} 1570 1571def Vector_StoreOp : Vector_Op<"store"> { 1572 let summary = "writes an n-D vector to an n-D slice of memory"; 1573 let description = [{ 1574 The 'vector.store' operation writes an n-D vector to an n-D slice of memory. 1575 It takes the vector value to be stored, a 'base' memref and an index for 1576 each memref dimension. The 'base' memref and indices determine the start 1577 memory address from which to write. Each index provides an offset for each 1578 memref dimension based on the element type of the memref. The shape of the 1579 vector value to store determines the shape of the slice written from the 1580 start memory address. The elements along each dimension of the slice are 1581 strided by the memref strides. Only unit strides are allowed along the most 1582 minor memref dimension. These constraints guarantee that elements written 1583 along the first dimension of the slice are contiguous in memory. 1584 1585 The memref element type can be a scalar or a vector type. If the memref 1586 element type is a scalar, it should match the element type of the value 1587 to store. If the memref element type is vector, it should match the type 1588 of the value to store. 1589 1590 Example 1: 1-D vector store on a scalar memref. 1591 ```mlir 1592 vector.store %valueToStore, %memref[%i, %j] : memref<200x100xf32>, vector<8xf32> 1593 ``` 1594 1595 Example 2: 1-D vector store on a vector memref. 1596 ```mlir 1597 vector.store %valueToStore, %memref[%i, %j] : memref<200x100xvector<8xf32>>, vector<8xf32> 1598 ``` 1599 1600 Example 3: 2-D vector store on a scalar memref. 1601 ```mlir 1602 vector.store %valueToStore, %memref[%i, %j] : memref<200x100xf32>, vector<4x8xf32> 1603 ``` 1604 1605 Example 4: 2-D vector store on a vector memref. 1606 ```mlir 1607 vector.store %valueToStore, %memref[%i, %j] : memref<200x100xvector<4x8xf32>>, vector<4x8xf32> 1608 ``` 1609 1610 Representation-wise, the 'vector.store' operation permits out-of-bounds 1611 writes. Support and implementation of out-of-bounds vector stores are 1612 target-specific. No assumptions should be made on the memory written out of 1613 bounds. Not all targets may support out-of-bounds vector stores. 1614 1615 Example 5: Potential out-of-bounds vector store. 1616 ```mlir 1617 vector.store %valueToStore, %memref[%index] : memref<?xf32>, vector<8xf32> 1618 ``` 1619 1620 Example 6: Explicit out-of-bounds vector store. 1621 ```mlir 1622 vector.store %valueToStore, %memref[%c0] : memref<7xf32>, vector<8xf32> 1623 ``` 1624 }]; 1625 1626 let arguments = (ins AnyVector:$valueToStore, 1627 Arg<AnyMemRef, "the reference to store to", 1628 [MemWrite]>:$base, 1629 Variadic<Index>:$indices); 1630 1631 let extraClassDeclaration = [{ 1632 MemRefType getMemRefType() { 1633 return getBase().getType().cast<MemRefType>(); 1634 } 1635 1636 VectorType getVectorType() { 1637 return getValueToStore().getType().cast<VectorType>(); 1638 } 1639 }]; 1640 1641 let hasFolder = 1; 1642 let hasVerifier = 1; 1643 1644 let assemblyFormat = "$valueToStore `,` $base `[` $indices `]` attr-dict " 1645 "`:` type($base) `,` type($valueToStore)"; 1646} 1647 1648def Vector_MaskedLoadOp : 1649 Vector_Op<"maskedload">, 1650 Arguments<(ins Arg<AnyMemRef, "", [MemRead]>:$base, 1651 Variadic<Index>:$indices, 1652 VectorOfRankAndType<[1], [I1]>:$mask, 1653 VectorOfRank<[1]>:$pass_thru)>, 1654 Results<(outs VectorOfRank<[1]>:$result)> { 1655 1656 let summary = "loads elements from memory into a vector as defined by a mask vector"; 1657 1658 let description = [{ 1659 The masked load reads elements from memory into a 1-D vector as defined 1660 by a base with indices and a 1-D mask vector. When the mask is set, the 1661 element is read from memory. Otherwise, the corresponding element is taken 1662 from a 1-D pass-through vector. Informally the semantics are: 1663 ``` 1664 result[0] := mask[0] ? base[i+0] : pass_thru[0] 1665 result[1] := mask[1] ? base[i+1] : pass_thru[1] 1666 etc. 1667 ``` 1668 The masked load can be used directly where applicable, or can be used 1669 during progressively lowering to bring other memory operations closer to 1670 hardware ISA support for a masked load. The semantics of the operation 1671 closely correspond to those of the `llvm.masked.load` 1672 [intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-load-intrinsics). 1673 1674 Examples: 1675 1676 ```mlir 1677 %0 = vector.maskedload %base[%i], %mask, %pass_thru 1678 : memref<?xf32>, vector<8xi1>, vector<8xf32> into vector<8xf32> 1679 1680 %1 = vector.maskedload %base[%i, %j], %mask, %pass_thru 1681 : memref<?x?xf32>, vector<16xi1>, vector<16xf32> into vector<16xf32> 1682 ``` 1683 }]; 1684 let extraClassDeclaration = [{ 1685 MemRefType getMemRefType() { 1686 return getBase().getType().cast<MemRefType>(); 1687 } 1688 VectorType getMaskVectorType() { 1689 return getMask().getType().cast<VectorType>(); 1690 } 1691 VectorType getPassThruVectorType() { 1692 return getPassThru().getType().cast<VectorType>(); 1693 } 1694 VectorType getVectorType() { 1695 return getResult().getType().cast<VectorType>(); 1696 } 1697 }]; 1698 let assemblyFormat = "$base `[` $indices `]` `,` $mask `,` $pass_thru attr-dict `:` " 1699 "type($base) `,` type($mask) `,` type($pass_thru) `into` type($result)"; 1700 let hasCanonicalizer = 1; 1701 let hasFolder = 1; 1702 let hasVerifier = 1; 1703} 1704 1705def Vector_MaskedStoreOp : 1706 Vector_Op<"maskedstore">, 1707 Arguments<(ins Arg<AnyMemRef, "", [MemWrite]>:$base, 1708 Variadic<Index>:$indices, 1709 VectorOfRankAndType<[1], [I1]>:$mask, 1710 VectorOfRank<[1]>:$valueToStore)> { 1711 1712 let summary = "stores elements from a vector into memory as defined by a mask vector"; 1713 1714 let description = [{ 1715 The masked store operation writes elements from a 1-D vector into memory 1716 as defined by a base with indices and a 1-D mask vector. When the mask is 1717 set, the corresponding element from the vector is written to memory. Otherwise, 1718 no action is taken for the element. Informally the semantics are: 1719 ``` 1720 if (mask[0]) base[i+0] = value[0] 1721 if (mask[1]) base[i+1] = value[1] 1722 etc. 1723 ``` 1724 The masked store can be used directly where applicable, or can be used 1725 during progressively lowering to bring other memory operations closer to 1726 hardware ISA support for a masked store. The semantics of the operation 1727 closely correspond to those of the `llvm.masked.store` 1728 [intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-store-intrinsics). 1729 1730 Examples: 1731 1732 ```mlir 1733 vector.maskedstore %base[%i], %mask, %value 1734 : memref<?xf32>, vector<8xi1>, vector<8xf32> 1735 1736 vector.maskedstore %base[%i, %j], %mask, %value 1737 : memref<?x?xf32>, vector<16xi1>, vector<16xf32> 1738 ``` 1739 }]; 1740 let extraClassDeclaration = [{ 1741 MemRefType getMemRefType() { 1742 return getBase().getType().cast<MemRefType>(); 1743 } 1744 VectorType getMaskVectorType() { 1745 return getMask().getType().cast<VectorType>(); 1746 } 1747 VectorType getVectorType() { 1748 return getValueToStore().getType().cast<VectorType>(); 1749 } 1750 }]; 1751 let assemblyFormat = 1752 "$base `[` $indices `]` `,` $mask `,` $valueToStore " 1753 "attr-dict `:` type($base) `,` type($mask) `,` type($valueToStore)"; 1754 let hasCanonicalizer = 1; 1755 let hasFolder = 1; 1756 let hasVerifier = 1; 1757} 1758 1759def Vector_GatherOp : 1760 Vector_Op<"gather">, 1761 Arguments<(ins Arg<AnyMemRef, "", [MemRead]>:$base, 1762 Variadic<Index>:$indices, 1763 VectorOfRankAndType<[1], [AnyInteger, Index]>:$index_vec, 1764 VectorOfRankAndType<[1], [I1]>:$mask, 1765 VectorOfRank<[1]>:$pass_thru)>, 1766 Results<(outs VectorOfRank<[1]>:$result)> { 1767 1768 let summary = "gathers elements from memory into a vector as defined by an index vector and mask"; 1769 1770 let description = [{ 1771 The gather operation gathers elements from memory into a 1-D vector as 1772 defined by a base with indices and an additional 1-D index vector, but 1773 only if the corresponding bit is set in a 1-D mask vector. Otherwise, the 1774 element is taken from a 1-D pass-through vector. Informally the semantics 1775 are: 1776 ``` 1777 result[0] := mask[0] ? base[index[0]] : pass_thru[0] 1778 result[1] := mask[1] ? base[index[1]] : pass_thru[1] 1779 etc. 1780 ``` 1781 The vector dialect leaves out-of-bounds behavior undefined. 1782 1783 The gather operation can be used directly where applicable, or can be used 1784 during progressively lowering to bring other memory operations closer to 1785 hardware ISA support for a gather. The semantics of the operation closely 1786 correspond to those of the `llvm.masked.gather` 1787 [intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-gather-intrinsics). 1788 1789 Examples: 1790 1791 ```mlir 1792 %0 = vector.gather %base[%c0][%v], %mask, %pass_thru 1793 : memref<?xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32> into vector<16xf32> 1794 1795 %1 = vector.gather %base[%i, %j][%v], %mask, %pass_thru 1796 : memref<16x16xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32> into vector<16xf32> 1797 ``` 1798 }]; 1799 let extraClassDeclaration = [{ 1800 MemRefType getMemRefType() { 1801 return getBase().getType().cast<MemRefType>(); 1802 } 1803 VectorType getIndexVectorType() { 1804 return getIndexVec().getType().cast<VectorType>(); 1805 } 1806 VectorType getMaskVectorType() { 1807 return getMask().getType().cast<VectorType>(); 1808 } 1809 VectorType getPassThruVectorType() { 1810 return getPassThru().getType().cast<VectorType>(); 1811 } 1812 VectorType getVectorType() { 1813 return getResult().getType().cast<VectorType>(); 1814 } 1815 }]; 1816 let assemblyFormat = 1817 "$base `[` $indices `]` `[` $index_vec `]` `,` " 1818 "$mask `,` $pass_thru attr-dict `:` type($base) `,` " 1819 "type($index_vec) `,` type($mask) `,` type($pass_thru) " 1820 "`into` type($result)"; 1821 let hasCanonicalizer = 1; 1822 let hasVerifier = 1; 1823} 1824 1825def Vector_ScatterOp : 1826 Vector_Op<"scatter">, 1827 Arguments<(ins Arg<AnyMemRef, "", [MemWrite]>:$base, 1828 Variadic<Index>:$indices, 1829 VectorOfRankAndType<[1], [AnyInteger, Index]>:$index_vec, 1830 VectorOfRankAndType<[1], [I1]>:$mask, 1831 VectorOfRank<[1]>:$valueToStore)> { 1832 1833 let summary = "scatters elements from a vector into memory as defined by an index vector and mask"; 1834 1835 let description = [{ 1836 The scatter operation scatters elements from a 1-D vector into memory as 1837 defined by a base with indices and an additional 1-D index vector, but 1838 only if the corresponding bit in a 1-D mask vector is set. Otherwise, no 1839 action is taken for that element. Informally the semantics are: 1840 ``` 1841 if (mask[0]) base[index[0]] = value[0] 1842 if (mask[1]) base[index[1]] = value[1] 1843 etc. 1844 ``` 1845 The vector dialect leaves out-of-bounds and repeated index behavior 1846 undefined. Underlying implementations may enforce strict sequential 1847 semantics for the latter, though. 1848 TODO: enforce the latter always? 1849 1850 The scatter operation can be used directly where applicable, or can be used 1851 during progressively lowering to bring other memory operations closer to 1852 hardware ISA support for a scatter. The semantics of the operation closely 1853 correspond to those of the `llvm.masked.scatter` 1854 [intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-scatter-intrinsics). 1855 1856 Examples: 1857 1858 ```mlir 1859 vector.scatter %base[%c0][%v], %mask, %value 1860 : memref<?xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32> 1861 1862 vector.scatter %base[%i, %j][%v], %mask, %value 1863 : memref<16x16xf32>, vector<16xi32>, vector<16xi1>, vector<16xf32> 1864 ``` 1865 }]; 1866 let extraClassDeclaration = [{ 1867 MemRefType getMemRefType() { 1868 return getBase().getType().cast<MemRefType>(); 1869 } 1870 VectorType getIndexVectorType() { 1871 return getIndexVec().getType().cast<VectorType>(); 1872 } 1873 VectorType getMaskVectorType() { 1874 return getMask().getType().cast<VectorType>(); 1875 } 1876 VectorType getVectorType() { 1877 return getValueToStore().getType().cast<VectorType>(); 1878 } 1879 }]; 1880 let assemblyFormat = 1881 "$base `[` $indices `]` `[` $index_vec `]` `,` " 1882 "$mask `,` $valueToStore attr-dict `:` type($base) `,` " 1883 "type($index_vec) `,` type($mask) `,` type($valueToStore)"; 1884 let hasCanonicalizer = 1; 1885 let hasVerifier = 1; 1886} 1887 1888def Vector_ExpandLoadOp : 1889 Vector_Op<"expandload">, 1890 Arguments<(ins Arg<AnyMemRef, "", [MemRead]>:$base, 1891 Variadic<Index>:$indices, 1892 VectorOfRankAndType<[1], [I1]>:$mask, 1893 VectorOfRank<[1]>:$pass_thru)>, 1894 Results<(outs VectorOfRank<[1]>:$result)> { 1895 1896 let summary = "reads elements from memory and spreads them into a vector as defined by a mask"; 1897 1898 let description = [{ 1899 The expand load reads elements from memory into a 1-D vector as defined 1900 by a base with indices and a 1-D mask vector. When the mask is set, the 1901 next element is read from memory. Otherwise, the corresponding element 1902 is taken from a 1-D pass-through vector. Informally the semantics are: 1903 ``` 1904 index = i 1905 result[0] := mask[0] ? base[index++] : pass_thru[0] 1906 result[1] := mask[1] ? base[index++] : pass_thru[1] 1907 etc. 1908 ``` 1909 Note that the index increment is done conditionally. 1910 1911 The expand load can be used directly where applicable, or can be used 1912 during progressively lowering to bring other memory operations closer to 1913 hardware ISA support for an expand. The semantics of the operation closely 1914 correspond to those of the `llvm.masked.expandload` 1915 [intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-expandload-intrinsics). 1916 1917 Examples: 1918 1919 ```mlir 1920 %0 = vector.expandload %base[%i], %mask, %pass_thru 1921 : memref<?xf32>, vector<8xi1>, vector<8xf32> into vector<8xf32> 1922 1923 %1 = vector.expandload %base[%i, %j], %mask, %pass_thru 1924 : memref<?x?xf32>, vector<16xi1>, vector<16xf32> into vector<16xf32> 1925 ``` 1926 }]; 1927 let extraClassDeclaration = [{ 1928 MemRefType getMemRefType() { 1929 return getBase().getType().cast<MemRefType>(); 1930 } 1931 VectorType getMaskVectorType() { 1932 return getMask().getType().cast<VectorType>(); 1933 } 1934 VectorType getPassThruVectorType() { 1935 return getPassThru().getType().cast<VectorType>(); 1936 } 1937 VectorType getVectorType() { 1938 return getResult().getType().cast<VectorType>(); 1939 } 1940 }]; 1941 let assemblyFormat = "$base `[` $indices `]` `,` $mask `,` $pass_thru attr-dict `:` " 1942 "type($base) `,` type($mask) `,` type($pass_thru) `into` type($result)"; 1943 let hasCanonicalizer = 1; 1944 let hasVerifier = 1; 1945} 1946 1947def Vector_CompressStoreOp : 1948 Vector_Op<"compressstore">, 1949 Arguments<(ins Arg<AnyMemRef, "", [MemWrite]>:$base, 1950 Variadic<Index>:$indices, 1951 VectorOfRankAndType<[1], [I1]>:$mask, 1952 VectorOfRank<[1]>:$valueToStore)> { 1953 1954 let summary = "writes elements selectively from a vector as defined by a mask"; 1955 1956 let description = [{ 1957 The compress store operation writes elements from a 1-D vector into memory 1958 as defined by a base with indices and a 1-D mask vector. When the mask is 1959 set, the corresponding element from the vector is written next to memory. 1960 Otherwise, no action is taken for the element. Informally the semantics are: 1961 ``` 1962 index = i 1963 if (mask[0]) base[index++] = value[0] 1964 if (mask[1]) base[index++] = value[1] 1965 etc. 1966 ``` 1967 Note that the index increment is done conditionally. 1968 1969 The compress store can be used directly where applicable, or can be used 1970 during progressively lowering to bring other memory operations closer to 1971 hardware ISA support for a compress. The semantics of the operation closely 1972 correspond to those of the `llvm.masked.compressstore` 1973 [intrinsic](https://llvm.org/docs/LangRef.html#llvm-masked-compressstore-intrinsics). 1974 1975 Examples: 1976 1977 ```mlir 1978 vector.compressstore %base[%i], %mask, %value 1979 : memref<?xf32>, vector<8xi1>, vector<8xf32> 1980 1981 vector.compressstore %base[%i, %j], %mask, %value 1982 : memref<?x?xf32>, vector<16xi1>, vector<16xf32> 1983 ``` 1984 }]; 1985 let extraClassDeclaration = [{ 1986 MemRefType getMemRefType() { 1987 return getBase().getType().cast<MemRefType>(); 1988 } 1989 VectorType getMaskVectorType() { 1990 return getMask().getType().cast<VectorType>(); 1991 } 1992 VectorType getVectorType() { 1993 return getValueToStore().getType().cast<VectorType>(); 1994 } 1995 }]; 1996 let assemblyFormat = 1997 "$base `[` $indices `]` `,` $mask `,` $valueToStore attr-dict `:` " 1998 "type($base) `,` type($mask) `,` type($valueToStore)"; 1999 let hasCanonicalizer = 1; 2000 let hasVerifier = 1; 2001} 2002 2003def Vector_ShapeCastOp : 2004 Vector_Op<"shape_cast", [NoSideEffect]>, 2005 Arguments<(ins AnyVector:$source)>, 2006 Results<(outs AnyVector:$result)> { 2007 let summary = "shape_cast casts between vector shapes"; 2008 let description = [{ 2009 The shape_cast operation casts between an n-D source vector shape and 2010 a k-D result vector shape (the element type remains the same). 2011 2012 If reducing rank (n > k), result dimension sizes must be a product 2013 of contiguous source dimension sizes. 2014 If expanding rank (n < k), source dimensions must factor into a 2015 contiguous sequence of destination dimension sizes. 2016 Each source dim is expanded (or contiguous sequence of source dims combined) 2017 in source dimension list order (i.e. 0 <= i < n), to produce a contiguous 2018 sequence of result dims (or a single result dim), in result dimension list 2019 order (i.e. 0 <= j < k). The product of all source dimension sizes and all 2020 result dimension sizes must match. 2021 2022 It is currently assumed that this operation does not require moving data, 2023 and that it will be folded away before lowering vector operations. 2024 2025 There is an exception to the folding expectation when targeting 2026 llvm.intr.matrix operations. We need a type conversion back and forth from a 2027 2-D MLIR vector to a 1-D flattened LLVM vector.shape_cast lowering to LLVM 2028 is supported in that particular case, for now. 2029 2030 Example: 2031 2032 ```mlir 2033 // Example casting to a lower vector rank. 2034 %1 = vector.shape_cast %0 : vector<5x1x4x3xf32> to vector<20x3xf32> 2035 2036 // Example casting to a higher vector rank. 2037 %3 = vector.shape_cast %2 : vector<10x12x8xf32> to vector<5x2x3x4x8xf32> 2038 2039 ``` 2040 }]; 2041 let extraClassDeclaration = [{ 2042 VectorType getSourceVectorType() { 2043 return getSource().getType().cast<VectorType>(); 2044 } 2045 VectorType getResultVectorType() { 2046 return getResult().getType().cast<VectorType>(); 2047 } 2048 }]; 2049 let assemblyFormat = "$source attr-dict `:` type($source) `to` type($result)"; 2050 let hasFolder = 1; 2051 let hasCanonicalizer = 1; 2052 let hasVerifier = 1; 2053} 2054 2055def Vector_BitCastOp : 2056 Vector_Op<"bitcast", [NoSideEffect, AllRanksMatch<["source", "result"]>]>, 2057 Arguments<(ins AnyVectorOfAnyRank:$source)>, 2058 Results<(outs AnyVectorOfAnyRank:$result)>{ 2059 let summary = "bitcast casts between vectors"; 2060 let description = [{ 2061 The bitcast operation casts between vectors of the same rank, the minor 1-D 2062 vector size is casted to a vector with a different element type but same 2063 bitwidth. In case of 0-D vectors, the bitwidth of element types must be 2064 equal. 2065 2066 Example: 2067 2068 ```mlir 2069 // Example casting to a smaller element type. 2070 %1 = vector.bitcast %0 : vector<5x1x4x3xf32> to vector<5x1x4x6xi16> 2071 2072 // Example casting to a bigger element type. 2073 %3 = vector.bitcast %2 : vector<10x12x8xi8> to vector<10x12x2xi32> 2074 2075 // Example casting to an element type of the same size. 2076 %5 = vector.bitcast %4 : vector<5x1x4x3xf32> to vector<5x1x4x3xi32> 2077 2078 // Example casting of 0-D vectors. 2079 %7 = vector.bitcast %6 : vector<f32> to vector<i32> 2080 ``` 2081 }]; 2082 let extraClassDeclaration = [{ 2083 VectorType getSourceVectorType() { 2084 return getSource().getType().cast<VectorType>(); 2085 } 2086 VectorType getResultVectorType() { 2087 return getResult().getType().cast<VectorType>(); 2088 } 2089 }]; 2090 let assemblyFormat = "$source attr-dict `:` type($source) `to` type($result)"; 2091 let hasFolder = 1; 2092 let hasVerifier = 1; 2093} 2094 2095def Vector_TypeCastOp : 2096 Vector_Op<"type_cast", [NoSideEffect, ViewLikeOpInterface]>, 2097 Arguments<(ins StaticShapeMemRefOf<[AnyType]>:$memref)>, 2098 Results<(outs AnyMemRef:$result)> { 2099 let summary = "type_cast op converts a scalar memref to a vector memref"; 2100 let description = [{ 2101 Performs a conversion from a memref with scalar element to a memref with a 2102 *single* vector element, copying the shape of the memref to the vector. This 2103 is the minimal viable operation that is required to makeke 2104 super-vectorization operational. It can be seen as a special case of the 2105 `view` operation but scoped in the super-vectorization context. 2106 2107 Syntax: 2108 2109 ``` 2110 operation ::= `vector.type_cast` ssa-use : memref-type to memref-type 2111 ``` 2112 2113 Example: 2114 2115 ```mlir 2116 %A = memref.alloc() : memref<5x4x3xf32> 2117 %VA = vector.type_cast %A : memref<5x4x3xf32> to memref<vector<5x4x3xf32>> 2118 ``` 2119 }]; 2120 2121 /// Build the canonical memRefType with a single vector. 2122 /// E.g. memref<4 x 5 x vector<6 x f32>> -> memref<vector<4 x 5 x 6 x f32>>. 2123 let builders = [OpBuilder<(ins "Value":$source)>]; 2124 2125 let extraClassDeclaration = [{ 2126 MemRefType getMemRefType() { 2127 return getMemref().getType().cast<MemRefType>(); 2128 } 2129 MemRefType getResultMemRefType() { 2130 return getResult().getType().cast<MemRefType>(); 2131 } 2132 // Implement ViewLikeOpInterface. 2133 Value getViewSource() { return getMemref(); } 2134 }]; 2135 2136 let assemblyFormat = [{ 2137 $memref attr-dict `:` type($memref) `to` type($result) 2138 }]; 2139 let hasVerifier = 1; 2140} 2141 2142def Vector_ConstantMaskOp : 2143 Vector_Op<"constant_mask", [NoSideEffect]>, 2144 Arguments<(ins I64ArrayAttr:$mask_dim_sizes)>, 2145 Results<(outs VectorOfAnyRankOf<[I1]>)> { 2146 let summary = "creates a constant vector mask"; 2147 let description = [{ 2148 Creates and returns a vector mask where elements of the result vector 2149 are set to '0' or '1', based on whether the element indices are contained 2150 within a hyper-rectangular region specified by the 'mask_dim_sizes' 2151 array attribute argument. Each element of the 'mask_dim_sizes' array, 2152 specifies an exclusive upper bound [0, mask-dim-size-element-value) 2153 for a unique dimension in the vector result. The conjunction of the ranges 2154 define a hyper-rectangular region within which elements values are set to 1 2155 (otherwise element values are set to 0). Each value of 'mask_dim_sizes' must 2156 be non-negative and not greater than the size of the corresponding vector 2157 dimension (as opposed to vector.create_mask which allows this). 2158 2159 Example: 2160 2161 ```mlir 2162 // create a constant vector mask of size 4x3xi1 with elements in range 2163 // 0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0). 2164 %1 = vector.constant_mask [3, 2] : vector<4x3xi1> 2165 2166 print %1 2167 columns 2168 0 1 2 2169 |------------ 2170 0 | 1 1 0 2171 rows 1 | 1 1 0 2172 2 | 1 1 0 2173 3 | 0 0 0 2174 ``` 2175 }]; 2176 2177 let extraClassDeclaration = [{ 2178 static StringRef getMaskDimSizesAttrStrName() { return "mask_dim_sizes"; } 2179 }]; 2180 let assemblyFormat = "$mask_dim_sizes attr-dict `:` type(results)"; 2181 let hasVerifier = 1; 2182} 2183 2184def Vector_CreateMaskOp : 2185 Vector_Op<"create_mask", [NoSideEffect]>, 2186 Arguments<(ins Variadic<Index>:$operands)>, 2187 Results<(outs VectorOfAnyRankOf<[I1]>)> { 2188 let summary = "creates a vector mask"; 2189 let description = [{ 2190 Creates and returns a vector mask where elements of the result vector 2191 are set to '0' or '1', based on whether the element indices are contained 2192 within a hyper-rectangular region specified by the operands. Specifically, 2193 each operand specifies a range [0, operand-value) for a unique dimension in 2194 the vector result. The conjunction of the operand ranges define a 2195 hyper-rectangular region within which elements values are set to 1 2196 (otherwise element values are set to 0). If operand-value is negative, it is 2197 treated as if it were zero, and if it is greater than the corresponding 2198 dimension size, it is treated as if it were equal to the dimension size. 2199 2200 Example: 2201 2202 ```mlir 2203 // create a vector mask of size 4x3xi1 where elements in range 2204 // 0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0). 2205 %1 = vector.create_mask %c3, %c2 : vector<4x3xi1> 2206 2207 print %1 2208 columns 2209 0 1 2 2210 |------------ 2211 0 | 1 1 0 2212 rows 1 | 1 1 0 2213 2 | 1 1 0 2214 3 | 0 0 0 2215 ``` 2216 }]; 2217 2218 let hasCanonicalizer = 1; 2219 let hasVerifier = 1; 2220 let assemblyFormat = "$operands attr-dict `:` type(results)"; 2221} 2222 2223def Vector_TransposeOp : 2224 Vector_Op<"transpose", [NoSideEffect, 2225 DeclareOpInterfaceMethods<VectorUnrollOpInterface, ["getShapeForUnroll"]>, 2226 PredOpTrait<"operand and result have same element type", 2227 TCresVTEtIsSameAsOpBase<0, 0>>]>, 2228 Arguments<(ins AnyVector:$vector, I64ArrayAttr:$transp)>, 2229 Results<(outs AnyVector:$result)> { 2230 let summary = "vector transpose operation"; 2231 let description = [{ 2232 Takes a n-D vector and returns the transposed n-D vector defined by 2233 the permutation of ranks in the n-sized integer array attribute. 2234 In the operation 2235 2236 ```mlir 2237 %1 = vector.transpose %0, [i_1, .., i_n] 2238 : vector<d_1 x .. x d_n x f32> 2239 to vector<d_trans[0] x .. x d_trans[n-1] x f32> 2240 ``` 2241 2242 the transp array [i_1, .., i_n] must be a permutation of [0, .., n-1]. 2243 2244 Example: 2245 2246 ```mlir 2247 %1 = vector.transpose %0, [1, 0] : vector<2x3xf32> to vector<3x2xf32> 2248 2249 [ [a, b, c], [ [a, d], 2250 [d, e, f] ] -> [b, e], 2251 [c, f] ] 2252 ``` 2253 }]; 2254 let builders = [ 2255 OpBuilder<(ins "Value":$vector, "ArrayRef<int64_t>":$transp)> 2256 ]; 2257 let extraClassDeclaration = [{ 2258 VectorType getVectorType() { 2259 return getVector().getType().cast<VectorType>(); 2260 } 2261 VectorType getResultType() { 2262 return getResult().getType().cast<VectorType>(); 2263 } 2264 void getTransp(SmallVectorImpl<int64_t> &results); 2265 static StringRef getTranspAttrStrName() { return "transp"; } 2266 }]; 2267 let assemblyFormat = [{ 2268 $vector `,` $transp attr-dict `:` type($vector) `to` type($result) 2269 }]; 2270 let hasCanonicalizer = 1; 2271 let hasFolder = 1; 2272 let hasVerifier = 1; 2273} 2274 2275def Vector_PrintOp : 2276 Vector_Op<"print", []>, Arguments<(ins AnyType:$source)> { 2277 let summary = "print operation (for testing and debugging)"; 2278 let description = [{ 2279 Prints the source vector (or scalar) to stdout in human readable 2280 format (for testing and debugging). No return value. 2281 2282 Example: 2283 2284 ```mlir 2285 %0 = arith.constant 0.0 : f32 2286 %1 = vector.broadcast %0 : f32 to vector<4xf32> 2287 vector.print %1 : vector<4xf32> 2288 2289 when lowered to LLVM, the vector print is unrolled into 2290 elementary printing method calls that at runtime will yield 2291 2292 ( 0.0, 0.0, 0.0, 0.0 ) 2293 2294 on stdout when linked with a small runtime support library, 2295 which only needs to provide a few printing methods (single 2296 value for all data types, opening/closing bracket, comma, 2297 newline). 2298 ``` 2299 }]; 2300 let extraClassDeclaration = [{ 2301 Type getPrintType() { 2302 return getSource().getType(); 2303 } 2304 }]; 2305 let assemblyFormat = "$source attr-dict `:` type($source)"; 2306} 2307 2308//===----------------------------------------------------------------------===// 2309// Ops used for supporting progressive lowering and conversion type changes. 2310// The Ops are typically not used directly by higher level dialects, but are 2311// used by intra-dialect rewriting rules to bring vector operations closer 2312// to the hardware ISA. 2313//===----------------------------------------------------------------------===// 2314 2315/// Vector dialect matrix multiplication op that operates on flattened 1-D 2316/// MLIR vectors. This is the counterpart of llvm.matrix.multiply in MLIR. 2317/// This may seem redundant with vector.contract but it serves the purposes of 2318/// more progressive lowering and localized type conversion on the path: 2319/// `vector<...x...xf32> -> vector<...xf32> -> !llvm<... x float>`. 2320def Vector_MatmulOp : Vector_Op<"matrix_multiply", [NoSideEffect, 2321 PredOpTrait<"lhs operand and result have same element type", 2322 TCresVTEtIsSameAsOpBase<0, 0>>, 2323 PredOpTrait<"rhs operand and result have same element type", 2324 TCresVTEtIsSameAsOpBase<0, 1>>]>, 2325 Arguments<( 2326 // TODO: tighten vector element types that make sense. 2327 ins VectorOfRankAndType<[1], 2328 [AnySignlessInteger, AnySignedInteger, Index, AnyFloat]>:$lhs, 2329 VectorOfRankAndType<[1], 2330 [AnySignlessInteger, AnySignedInteger, Index, AnyFloat]>:$rhs, 2331 I32Attr:$lhs_rows, I32Attr:$lhs_columns, I32Attr:$rhs_columns)>, 2332 Results<( 2333 outs VectorOfRankAndType<[1], 2334 [AnySignlessInteger, AnySignedInteger, Index, AnyFloat]>:$res)> 2335{ 2336 let summary = "Vector matrix multiplication op that operates on flattened 1-D" 2337 " MLIR vectors"; 2338 let description = [{ 2339 This is the counterpart of llvm.matrix.multiply in MLIR. It serves the 2340 purposes of more progressive lowering and localized type conversion. 2341 Higher levels typically lower matrix multiplications into 'vector.contract' 2342 operations. Subsequent rewriting rule progressively lower these operations 2343 into 'vector.matrix_multiply' operations to bring the operations closer 2344 to the hardware ISA. 2345 2346 The ‘vector.matrix_multiply’ op treats `lhs` as matrix with <lhs_rows> rows 2347 and <lhs_columns> columns, `rhs` as matrix with <lhs_columns> rows and 2348 <rhs_columns> and multiplies them. The result matrix is returned embedded in 2349 the result vector. 2350 2351 Also see: 2352 2353 http://llvm.org/docs/LangRef.html#llvm-matrix-multiply-intrinsic 2354 2355 Example: 2356 2357 ```mlir 2358 %C = vector.matrix_multiply %A, %B 2359 { lhs_rows = 4: i32, lhs_columns = 16: i32 , rhs_columns = 3: i32 } : 2360 (vector<64xf64>, vector<48xf64>) -> vector<12xf64> 2361 ``` 2362 }]; 2363 let builders = [ 2364 OpBuilder<(ins "Value":$lhs, "Value":$rhs, "unsigned":$lhsRows, 2365 "unsigned":$lhsColumns, "unsigned":$rhsColumns), 2366 [{ 2367 $_state.addOperands({lhs, rhs}); 2368 $_state.addAttribute("lhs_rows",$_builder.getI32IntegerAttr(lhsRows)); 2369 $_state.addAttribute("lhs_columns",$_builder.getI32IntegerAttr(lhsColumns)); 2370 $_state.addAttribute("rhs_columns",$_builder.getI32IntegerAttr(rhsColumns)); 2371 $_state.addTypes(VectorType::get(lhsRows * rhsColumns, 2372 lhs.getType().cast<VectorType>().getElementType())); 2373 }]>, 2374 ]; 2375 let assemblyFormat = "$lhs `,` $rhs attr-dict " 2376 "`:` `(` type($lhs) `,` type($rhs) `)` `->` type($res)"; 2377} 2378 2379/// Vector dialect matrix tranposition op that operates on flattened 1-D 2380/// MLIR vectors. This is the counterpart of llvm.matrix.transpose in MLIR. 2381/// This may seem redundant with vector.transpose but it serves the purposes of 2382/// more progressive lowering and localized type conversion on the path: 2383/// `vector<...x...xf32> -> vector<...xf32> -> !llvm<... x float>`. 2384def Vector_FlatTransposeOp : Vector_Op<"flat_transpose", [NoSideEffect, 2385 PredOpTrait<"source operand and result have same element type", 2386 TCresVTEtIsSameAsOpBase<0, 0>>]>, 2387 Arguments<( 2388 // TODO: tighten vector element types that make sense. 2389 ins VectorOfRankAndType<[1], 2390 [AnySignlessInteger, AnySignedInteger, Index, AnyFloat]>:$matrix, 2391 I32Attr:$rows, I32Attr:$columns)>, 2392 Results<( 2393 outs VectorOfRankAndType<[1], 2394 [AnySignlessInteger, AnySignedInteger, Index, AnyFloat]>:$res)> { 2395 let summary = "Vector matrix transposition on flattened 1-D MLIR vectors"; 2396 let description = [{ 2397 This is the counterpart of llvm.matrix.transpose in MLIR. It serves 2398 the purposes of more progressive lowering and localized type conversion. 2399 Higher levels typically lower matrix tranpositions into 'vector.transpose' 2400 operations. Subsequent rewriting rule progressively lower these operations 2401 into 'vector.flat_transpose' operations to bring the operations closer 2402 to the hardware ISA. 2403 2404 The ‘vector.flat_transpose’ op treats the 1-D input `matrix` as 2405 a 2-D matrix with <rows> rows and <columns> columns, and returns the 2406 transposed matrix in flattened form in 'res'. 2407 2408 Also see: 2409 2410 http://llvm.org/docs/LangRef.html#llvm-matrix-transpose-intrinsic 2411 2412 Example: 2413 2414 ```mlir 2415 %1 = vector.flat_transpose %0 { rows = 4: i32, columns = 4: i32 } 2416 : (vector<16xf32>) -> vector<16xf32> 2417 ``` 2418 }]; 2419 let assemblyFormat = "$matrix attr-dict `:` type($matrix) `->` type($res)"; 2420} 2421 2422//===----------------------------------------------------------------------===// 2423// SplatOp 2424//===----------------------------------------------------------------------===// 2425 2426def Vector_SplatOp : Vector_Op<"splat", [ 2427 NoSideEffect, 2428 TypesMatchWith<"operand type matches element type of result", 2429 "aggregate", "input", 2430 "$_self.cast<VectorType>().getElementType()"> 2431 ]> { 2432 let summary = "vector splat or broadcast operation"; 2433 let description = [{ 2434 Broadcast the operand to all elements of the result vector. The operand is 2435 required to be of integer/index/float type. 2436 2437 Example: 2438 2439 ```mlir 2440 %s = arith.constant 10.1 : f32 2441 %t = vector.splat %s : vector<8x16xi32> 2442 ``` 2443 }]; 2444 2445 let arguments = (ins AnyTypeOf<[AnySignlessInteger, Index, AnyFloat], 2446 "integer/index/float type">:$input); 2447 let results = (outs AnyVectorOfAnyRank:$aggregate); 2448 2449 let builders = [ 2450 OpBuilder<(ins "Value":$element, "Type":$aggregateType), 2451 [{ build($_builder, $_state, aggregateType, element); }]>]; 2452 let assemblyFormat = "$input attr-dict `:` type($aggregate)"; 2453 2454 let hasFolder = 1; 2455} 2456 2457//===----------------------------------------------------------------------===// 2458// VectorScaleOp 2459//===----------------------------------------------------------------------===// 2460 2461// TODO: In the future, we might want to have scalable vectors with different 2462// scales for different dimensions. E.g.: vector<[16]x[16]xf32>, in 2463// which case we might need to add an index to 'vscale' to select one 2464// of them. In order to support GPUs, we might also want to differentiate 2465// between a 'global' scale, a scale that's fixed throughout the 2466// execution, and a 'local' scale that is fixed but might vary with each 2467// call to the function. For that, it might be useful to have a 2468// 'vector.scale.global' and a 'vector.scale.local' operation. 2469def VectorScaleOp : Vector_Op<"vscale", 2470 [NoSideEffect]> { 2471 let summary = "Load vector scale size"; 2472 let description = [{ 2473 The `vscale` op returns the scale of the scalable vectors, a positive 2474 integer value that is constant at runtime but unknown at compile-time. 2475 The scale of the vector indicates the multiplicity of the vectors and 2476 vector operations. For example, a `vector<[4]xi32>` is equivalent to 2477 `vscale` consecutive `vector<4xi32>`; and an operation on a 2478 `vector<[4]xi32>` is equivalent to performing that operation `vscale` 2479 times, once on each `<4xi32>` segment of the scalable vector. The `vscale` 2480 op can be used to calculate the step in vector-length agnostic (VLA) loops. 2481 Right now we only support one contiguous set of scalable dimensions, all of 2482 them grouped and scaled with the value returned by 'vscale'. 2483 }]; 2484 let results = (outs Index:$res); 2485 let assemblyFormat = "attr-dict"; 2486} 2487 2488//===----------------------------------------------------------------------===// 2489// VectorScanOp 2490//===----------------------------------------------------------------------===// 2491 2492def Vector_ScanOp : 2493 Vector_Op<"scan", [NoSideEffect, 2494 AllTypesMatch<["source", "dest"]>, 2495 AllTypesMatch<["initial_value", "accumulated_value"]> ]>, 2496 Arguments<(ins Vector_CombiningKindAttr:$kind, 2497 AnyVector:$source, 2498 AnyVectorOfAnyRank:$initial_value, 2499 I64Attr:$reduction_dim, 2500 BoolAttr:$inclusive)>, 2501 Results<(outs AnyVector:$dest, 2502 AnyVectorOfAnyRank:$accumulated_value)> { 2503 let summary = "Scan operation"; 2504 let description = [{ 2505 Performs an inclusive/exclusive scan on an n-D vector along a single 2506 dimension returning an n-D result vector using the given 2507 operation (add/mul/min/max for int/fp and and/or/xor for 2508 int only) and a specified value for the initial value. The operator 2509 returns the result of scan as well as the result of the last 2510 reduction in the scan. 2511 2512 Example: 2513 2514 ```mlir 2515 %1:2 = vector.scan <add>, %0, %acc {inclusive = false, reduction_dim = 1 : i64} : 2516 vector<4x8x16x32xf32>, vector<4x16x32xf32> 2517 ``` 2518 }]; 2519 let builders = [ 2520 OpBuilder<(ins "Value":$source, "Value":$initial_value, 2521 "CombiningKind":$kind, 2522 CArg<"int64_t", "0">:$reduction_dim, 2523 CArg<"bool", "true">:$inclusive)> 2524 ]; 2525 let extraClassDeclaration = [{ 2526 static StringRef getKindAttrStrName() { return "kind"; } 2527 static StringRef getReductionDimAttrStrName() { return "reduction_dim"; } 2528 VectorType getSourceType() { 2529 return getSource().getType().cast<VectorType>(); 2530 } 2531 VectorType getDestType() { 2532 return getDest().getType().cast<VectorType>(); 2533 } 2534 VectorType getAccumulatorType() { 2535 return getAccumulatedValue().getType().cast<VectorType>(); 2536 } 2537 VectorType getInitialValueType() { 2538 return getInitialValue().getType().cast<VectorType>(); 2539 } 2540 }]; 2541 let assemblyFormat = 2542 "$kind `,` $source `,` $initial_value attr-dict `:` " 2543 "type($source) `,` type($initial_value) "; 2544 let hasVerifier = 1; 2545} 2546 2547def Vector_YieldOp : Vector_Op<"yield", [ 2548 NoSideEffect, ReturnLike, Terminator]> { 2549 let summary = "Terminates and yields values from vector regions."; 2550 let description = [{ 2551 "vector.yield" yields an SSA value from the Vector dialect op region and 2552 terminates the regions. The semantics of how the values are yielded is 2553 defined by the parent operation. 2554 If "vector.yield" has any operands, the operands must correspond to the 2555 parent operation's results. 2556 If the parent operation defines no value the vector.yield may be omitted 2557 when printing the region. 2558 }]; 2559 2560 let arguments = (ins Variadic<AnyType>:$operands); 2561 2562 let builders = [ 2563 OpBuilder<(ins), [{ /* nothing to do */ }]>, 2564 ]; 2565 2566 let assemblyFormat = "attr-dict ($operands^ `:` type($operands))?"; 2567} 2568 2569def Vector_WarpExecuteOnLane0Op : Vector_Op<"warp_execute_on_lane_0", 2570 [DeclareOpInterfaceMethods<RegionBranchOpInterface, ["areTypesCompatible"]>, 2571 SingleBlockImplicitTerminator<"vector::YieldOp">, 2572 RecursiveSideEffects]> { 2573 let summary = "Executes operations in the associated region on thread #0 of a" 2574 "SPMD program"; 2575 let description = [{ 2576 `warp_execute_on_lane_0` is an operation used to bridge the gap between 2577 vector programming and SPMD programming model like GPU SIMT. It allows to 2578 trivially convert a region of vector code meant to run on a multiple threads 2579 into a valid SPMD region and then allows incremental transformation to 2580 distribute vector operations on the threads. 2581 2582 Any code present in the region would only be executed on first thread/lane 2583 based on the `laneid` operand. The `laneid` operand is an integer ID between 2584 [0, `warp_size`). The `warp_size` attribute indicates the number of lanes in 2585 a warp. 2586 2587 Operands are vector values distributed on all lanes that may be used by 2588 the single lane execution. The matching region argument is a vector of all 2589 the values of those lanes available to the single active lane. The 2590 distributed dimension is implicit based on the shape of the operand and 2591 argument. the properties of the distribution may be described by extra 2592 attributes (e.g. affine map). 2593 2594 Return values are distributed on all lanes using laneId as index. The 2595 vector is distributed based on the shape ratio between the vector type of 2596 the yield and the result type. 2597 If the shapes are the same this means the value is broadcasted to all lanes. 2598 In the future the distribution can be made more explicit using affine_maps 2599 and will support having multiple Ids. 2600 2601 Therefore the `warp_execute_on_lane_0` operations allow to implicitly copy 2602 between lane0 and the lanes of the warp. When distributing a vector 2603 from lane0 to all the lanes, the data are distributed in a block cyclic way. 2604 For exemple `vector<64xf32>` gets distributed on 32 threads and map to 2605 `vector<2xf32>` where thread 0 contains vector[0] and vector[1]. 2606 2607 During lowering values passed as operands and return value need to be 2608 visible to different lanes within the warp. This would usually be done by 2609 going through memory. 2610 2611 The region is *not* isolated from above. For values coming from the parent 2612 region not going through operands only the lane 0 value will be accesible so 2613 it generally only make sense for uniform values. 2614 2615 Example: 2616 ``` 2617 // Execute in parallel on all threads/lanes. 2618 vector.warp_execute_on_lane_0 (%laneid)[32] { 2619 // Serial code running only on thread/lane 0. 2620 ... 2621 } 2622 // Execute in parallel on all threads/lanes. 2623 ``` 2624 2625 This may be lowered to an scf.if region as below: 2626 ``` 2627 // Execute in parallel on all threads/lanes. 2628 %cnd = arith.cmpi eq, %laneid, %c0 : index 2629 scf.if %cnd { 2630 // Serial code running only on thread/lane 0. 2631 ... 2632 } 2633 // Execute in parallel on all threads/lanes. 2634 ``` 2635 2636 When the region has operands and/or return values: 2637 ``` 2638 // Execute in parallel on all threads/lanes. 2639 %0 = vector.warp_execute_on_lane_0(%laneid)[32] 2640 args(%v0 : vector<4xi32>) -> (vector<1xf32>) { 2641 ^bb0(%arg0 : vector<128xi32>) : 2642 // Serial code running only on thread/lane 0. 2643 ... 2644 vector.yield %1 : vector<32xf32> 2645 } 2646 // Execute in parallel on all threads/lanes. 2647 ``` 2648 2649 values at the region boundary would go through memory: 2650 ``` 2651 // Execute in parallel on all threads/lanes. 2652 ... 2653 // Store the data from each thread into memory and Synchronization. 2654 %tmp0 = memreg.alloc() : memref<128xf32> 2655 %tmp1 = memreg.alloc() : memref<32xf32> 2656 %cnd = arith.cmpi eq, %laneid, %c0 : index 2657 vector.store %v0, %tmp0[%laneid] : memref<128xf32>, vector<4xf32> 2658 some_synchronization_primitive 2659 scf.if %cnd { 2660 // Serialized code running only on thread 0. 2661 // Load the data from all the threads into a register from thread 0. This 2662 // allow threads 0 to access data from all the threads. 2663 %arg0 = vector.load %tmp0[%c0] : memref<128xf32>, vector<128xf32> 2664 ... 2665 // Store the data from thread 0 into memory. 2666 vector.store %1, %tmp1[%c0] : memref<32xf32>, vector<32xf32> 2667 } 2668 // Synchronization and load the data in a block cyclic way so that the 2669 // vector is distributed on all threads. 2670 some_synchronization_primitive 2671 %0 = vector.load %tmp1[%laneid] : memref<32xf32>, vector<32xf32> 2672 // Execute in parallel on all threads/lanes. 2673 ``` 2674 2675 }]; 2676 2677 let hasVerifier = 1; 2678 let hasCustomAssemblyFormat = 1; 2679 let arguments = (ins Index:$laneid, I64Attr:$warp_size, 2680 Variadic<AnyType>:$args); 2681 let results = (outs Variadic<AnyType>:$results); 2682 let regions = (region SizedRegion<1>:$warpRegion); 2683 2684 let skipDefaultBuilders = 1; 2685 let builders = [ 2686 OpBuilder<(ins "Value":$laneid, "int64_t":$warpSize)>, 2687 OpBuilder<(ins "TypeRange":$resultTypes, "Value":$laneid, 2688 "int64_t":$warpSize)>, 2689 // `blockArgTypes` are different than `args` types as they are they 2690 // represent all the `args` instances visibile to lane 0. Therefore we need 2691 // to explicit pass the type. 2692 OpBuilder<(ins "TypeRange":$resultTypes, "Value":$laneid, 2693 "int64_t":$warpSize, "ValueRange":$args, 2694 "TypeRange":$blockArgTypes)> 2695 ]; 2696 2697 let extraClassDeclaration = [{ 2698 bool isDefinedOutsideOfRegion(Value value) { 2699 return !getRegion().isAncestor(value.getParentRegion()); 2700 } 2701 }]; 2702} 2703 2704#endif // VECTOR_OPS 2705