施工実績

The duration of the fresh assortment need certainly to equal this new score of your own vector form of. Aired size must always enter-bounds. Its lack of it elective in_bounds characteristic shows that one aspect of your own import (with the exception of broadcasts) tends to be out-of-bounds. A good vector.transfer_see can be lowered so you can an easy load if all size is specified to get contained in this bounds no mask try given.

Keep in mind that inside_bounds was given to possess effect size and not enter in proportions. Brand new first step of transfer, we.e., %A[%expr1, %expr2, %expr3, %expr4] from the analogy less than, is expected to settle-bounds and also as indices try growing, accesses will get go out-of-bounds.

So it procedure is called ‘see because of the resistance in order to ‘weight as very-vector granularity may be perhaps not representable with an individual knowledge register. A great vector.transfer_see try therefore a middle-level abstraction one helps very-vectorization having non-effecting padding getting full-tile merely operations.

This procedure usually checks out a slice performing within %A[%expr1, %expr2, %expr3, %expr4] . How big the fresh cut was step three along d2 and you may 5 together d0, and so the cut is actually: %A[%expr1 : %expr1 + 5, %expr2, %expr3:%expr3 + 3, %expr4]

That slice needs to be read into a vector<3x4x5xf32> . Since the permutation map is not full rank, there must be a broadcast along vector dimension 1 .

Toward a good GPU one can next chart we , j , k to help you prevents and you will posts. See that the fresh new short term shops footprint try step 3 * 5 viewpoints however, step three * 4 * 5 beliefs happen to be transferred between %A beneficial and you can %tmp .

in which broadcast shows away from feature 0 to anyone else over the specified measurement. Now, brand new brief shop footprint try 3 * 5 values the same amount of investigation as step 3 * 5 viewpoints directed. An extra step one aired needs. On a good GPU which shown could be then followed using a warp-shuffle in the event the loop j was basically mapped to threadIdx.x .

vector.transfer_build (::mlir::vector::TransferWriteOp) ¶

The brand new vector.transfer_develop op performs a write from a beneficial vector, given as the earliest operand, on a piece within a good MemRef or a placed Tensor regarding the same foot elemental type of, offered as its next operand.

A vector memref/tensor operand must have its vector element type match a suffix (shape and element type) of the vector (e.g. memref<3x2x6x4x3xf32>, vector<1x1x4x3xf32>). If the operand is a tensor, the operation returns a new tensor of the same type.

Operands: ¶

New cut is actually after that defined of the a complete-rating directory when you look at the MemRef/Tensor, supplied since operands [2 .. dos + rank(memref/tensor)) .

New permutation_map trait is actually a keen affine-map which determine the brand new transposition with the slice to fit new vector figure. The newest permutation chart could be implicit and omitted off parsing and you can printing if it is the fresh canonical small identity map (we.age. if this will not permute any aspect). Weighed against transfer_see , generate ops try not to possess aired proportions.

A recommended SSA worth mask of the identical profile as the vector particular tends to be specified in order to mask out issue. Factors whoever associated hide feature is actually 0 is masked aside.

The absence of that guyspy it inside_bounds trait implies that people dimensions of your import could be out-of-bounds. A beneficial vector.transfer_create is going to be paid off to a simple shop if all of the proportions try given to-be in this bounds no hide was given.

The length of the assortment need certainly to equal brand new score of your own vector sort of. The absence of this optional inside_bounds attribute shows that any dimensions of transfer can be out-of-bounds. Good vector.transfer_write will be paid off so you’re able to a straightforward store in the event that all the dimensions is given to-be contained in this bounds without mask was specified.