{-# OPTIONS_GHC -Wno-unused-imports #-}
{-# OPTIONS_HADDOCK hide, prune, not-home #-}

module MLIR.AST.Dialect.Generated.Vector where

import Prelude (Int, Double, Maybe(..), Bool(..), (++), (<$>), ($), (<>), Show)
import qualified Prelude
import Data.Int (Int64)
import qualified Data.Maybe
import Data.Array (Ix)
import qualified Data.Array.IArray as IArray
import qualified Data.ByteString as BS
import qualified Data.Map.Strict as M
import qualified Control.Monad

import MLIR.AST ( Attribute(..), Type(..), AbstractOperation(..), ResultTypes(..)
                , Location(..), Signedness(..), DenseElements(..)
                , NamedAttributes, Name
                , pattern NoAttrs )
import qualified MLIR.AST as AST
import MLIR.AST.Builder (Value, EndOfBlock, MonadBlockBuilder, RegionBuilderT)
import qualified MLIR.AST.Builder as AST
import qualified MLIR.AST.IStorableArray as AST
import qualified MLIR.AST.PatternUtil as PatternUtil
import qualified MLIR.AST.Dialect.Affine as Affine

-- * vscale
-- $vscale
-- 
-- The @vscale@ op returns the scale of the scalable vectors, a positive
-- integer value that is constant at runtime but unknown at compile-time.
-- The scale of the vector indicates the multiplicity of the vectors and
-- vector operations. For example, a @vector\<[4]xi32>@ is equivalent to
-- @vscale@ consecutive @vector\<4xi32>@; and an operation on a
-- @vector\<[4]xi32>@ is equivalent to performing that operation @vscale@
-- times, once on each @\<4xi32>@ segment of the scalable vector. The @vscale@
-- op can be used to calculate the step in vector-length agnostic (VLA) loops.
-- Right now we only support one contiguous set of scalable dimensions, all of
-- them grouped and scaled with the value returned by \'vscale\'.
--   

-- | A builder for @vector.vscale@.
vscale :: () => MonadBlockBuilder m => Type -> m Value
vscale :: Type -> m Value
vscale Type
ty0    = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.vscale"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = []
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * bitcast
-- $bitcast
-- 
-- The bitcast operation casts between vectors of the same rank, the minor 1-D
-- vector size is casted to a vector with a different element type but same
-- bitwidth. In case of 0-D vectors, the bitwidth of element types must be
-- equal.
-- 
-- Example:
-- 
-- @
-- // Example casting to a smaller element type.
-- %1 = vector.bitcast %0 : vector\<5x1x4x3xf32> to vector\<5x1x4x6xi16>
-- 
-- // Example casting to a bigger element type.
-- %3 = vector.bitcast %2 : vector\<10x12x8xi8> to vector\<10x12x2xi32>
-- 
-- // Example casting to an element type of the same size.
-- %5 = vector.bitcast %4 : vector\<5x1x4x3xf32> to vector\<5x1x4x3xi32>
-- 
-- // Example casting of 0-D vectors.
-- %7 = vector.bitcast %6 : vector\<f32> to vector\<i32>
-- @
--   

-- | A pattern for @vector.bitcast@.
pattern BitCast :: () => () => Location -> Type -> operand -> AbstractOperation operand
pattern $bBitCast :: Location -> Type -> operand -> AbstractOperation operand
$mBitCast :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> r) -> (Void# -> r) -> r
BitCast loc ty0 source_  = Operation
          { opName = "vector.bitcast"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [source_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.bitcast@.
bitcast :: () => MonadBlockBuilder m => Type -> Value -> m Value
bitcast :: Type -> Value -> m Value
bitcast Type
ty0 Value
source_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.bitcast"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
source_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * broadcast
-- $broadcast
-- 
-- Broadcasts the scalar or k-D vector value in the source operand
-- to a n-D result vector such that the broadcast makes sense, i.e.,
-- the source operand is duplicated to match the given rank and sizes
-- in the result vector. The legality rules are:
-- * the source operand must have the same element type as the result type
-- * a k-D vector \<s_1 x .. x s_k x type> can be broadcast to
--   a n-D vector \<t_1 x .. x t_n x type> if
--    * k \<= n, and
--    * the sizes in the trailing dimensions n-k \< i \<= n with j=i+k-n
--       match exactly as s_j = t_i or s_j = 1:
--    @
--        t_1 x   ..  t_n-k x t_n-k+1 x .. x t_i x .. x t_n
--                            s_1     x .. x s_j x .. x s_k
--            \<duplication>         \<potential stretch>
--    @
--    * in addition, any scalable unit dimension, @[1]@, must match exactly.
-- 
-- The source operand is duplicated over all the missing leading dimensions
-- and stretched over the trailing dimensions where the source has a non-equal
-- dimension of 1. These rules imply that any scalar broadcast (k=0) to any
-- shaped vector with the same element type is always legal.
-- 
-- Example:
-- 
-- @
-- %0 = arith.constant 0.0 : f32
-- %1 = vector.broadcast %0 : f32 to vector\<16xf32>
-- %2 = vector.broadcast %1 : vector\<16xf32> to vector\<4x16xf32>
-- @
--   

-- | A pattern for @vector.broadcast@.
pattern Broadcast :: () => () => Location -> Type -> operand -> AbstractOperation operand
pattern $bBroadcast :: Location -> Type -> operand -> AbstractOperation operand
$mBroadcast :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> r) -> (Void# -> r) -> r
Broadcast loc ty0 source_  = Operation
          { opName = "vector.broadcast"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [source_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.broadcast@.
broadcast :: () => MonadBlockBuilder m => Type -> Value -> m Value
broadcast :: Type -> Value -> m Value
broadcast Type
ty0 Value
source_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.broadcast"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
source_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * compressstore
-- $compressstore
-- 
-- The compress store operation writes elements from a vector into memory as
-- defined by a base with indices and a mask vector. Compression only applies
-- to the innermost dimension. When the mask is set, the corresponding element
-- from the vector is written next to memory.  Otherwise, no action is taken
-- for the element. Informally the semantics are:
-- 
-- @
-- index = i
-- if (mask[0]) base[index++] = value[0]
-- if (mask[1]) base[index++] = value[1]
-- etc.
-- @
-- 
-- Note that the index increment is done conditionally.
-- 
-- If a mask bit is set and the corresponding index is out-of-bounds for the
-- given base, the behavior is undefined. If a mask bit is not set, no value
-- is stored regardless of the index, and the index is allowed to be
-- out-of-bounds.
-- 
-- The compress store can be used directly where applicable, or can be used
-- during progressively lowering to bring other memory operations closer to
-- hardware ISA support for a compress. The semantics of the operation closely
-- correspond to those of the @llvm.masked.compressstore@
-- intrinsic.
-- 
-- Note, at the moment this Op is only available for fixed-width vectors.
-- 
-- Examples:
-- 
-- @
-- vector.compressstore %base[%i], %mask, %value
--   : memref\<?xf32>, vector\<8xi1>, vector\<8xf32>
-- 
-- vector.compressstore %base[%i, %j], %mask, %value
--   : memref\<?x?xf32>, vector\<16xi1>, vector\<16xf32>
-- @
--   

-- | A builder for @vector.compressstore@.
compressstore :: () => MonadBlockBuilder m => Value -> [Value] -> Value -> Value -> m ()
compressstore :: Value -> [Value] -> Value -> Value -> m ()
compressstore  Value
base_ [Value]
indices_ Value
mask_ Value
valueToStore_   = do
  m [Value] -> m ()
forall (f :: * -> *) a. Functor f => f a -> f ()
Control.Monad.void (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.compressstore"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit []
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_) [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
mask_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
valueToStore_)])
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * constant_mask
-- $constant_mask
-- 
-- Creates and returns a vector mask where elements of the result vector
-- are set to \'0\' or \'1\', based on whether the element indices are contained
-- within a hyper-rectangular region specified by the \'mask_dim_sizes\'
-- array attribute argument. Each element of the \'mask_dim_sizes\' array,
-- specifies an exclusive upper bound [0, mask-dim-size-element-value)
-- for a unique dimension in the vector result. The conjunction of the ranges
-- define a hyper-rectangular region within which elements values are set to 1
-- (otherwise element values are set to 0). Each value of \'mask_dim_sizes\' must
-- be non-negative and not greater than the size of the corresponding vector
-- dimension (as opposed to vector.create_mask which allows this). Sizes that
-- correspond to scalable dimensions are implicitly multiplied by vscale,
-- though currently only zero (none set) or the size of the dim/vscale
-- (all set) are supported.
-- 
-- Example:
-- 
-- @
-- // create a constant vector mask of size 4x3xi1 with elements in range
-- // 0 \<= row \<= 2 and 0 \<= col \<= 1 are set to 1 (others to 0).
-- %1 = vector.constant_mask [3, 2] : vector\<4x3xi1>
-- 
-- print %1
--               columns
--             0    1    2
--           |------------
--         0 | 1    1    0
--   rows  1 | 1    1    0
--         2 | 1    1    0
--         3 | 0    0    0
-- @
--   

-- * contract
-- $contract
-- 
-- Computes the sum of products of vector elements along contracting
-- dimension pairs from 2 vectors of rank M and N respectively, adds this
-- intermediate result to the accumulator argument of rank K, and returns a
-- vector result of rank K (where K = num_lhs_free_dims + num_rhs_free_dims +
-- num_batch_dims (see dimension type descriptions below)). For K = 0 (no
-- free or batch dimensions), the accumulator and output are a scalar.
-- 
-- If operands and the result have types of different bitwidths, operands are
-- promoted to have the same bitwidth as the result before performing the
-- contraction. For integer types, only signless integer types are supported,
-- and the promotion happens via sign extension.
-- 
-- An iterator type attribute list must be specified, where each element of
-- the list represents an iterator with one of the following types:
-- 
-- *   \"reduction\": reduction dimensions are present in the lhs and rhs
--     arguments but not in the output (and accumulator
--     argument). These are the dimensions along which the vector
--     contraction op computes the sum of products, and
--     contracting dimension pair dimension sizes must match
--     between lhs/rhs.
-- 
-- *   \"parallel\": Batch dimensions are iterator type \"parallel\", and
--     are non-contracting dimensions present in the lhs, rhs and
--     output. The lhs/rhs co-iterate along the batch dimensions,
--     which should be expressed in their indexing maps.
-- 
--     Free dimensions are iterator type \"parallel\", and are
--     non-contraction, non-batch dimensions accessed by either the
--     lhs or rhs (but not both). The lhs and rhs free dimensions
--     are unrelated to each other and do not co-iterate, which
--     should be expressed in their indexing maps.
-- 
-- An indexing map attribute list must be specified with an entry for lhs, rhs
-- and acc arguments. An indexing map attribute specifies a mapping from each
-- iterator in the iterator type list, to each dimension of an N-D vector.
-- 
-- An optional kind attribute may be used to specify the combining function
-- between the intermediate result and accumulator argument of rank K. This
-- attribute can take the values @add@/@mul@/@minsi@/@minui@/@maxsi@/@maxui@
-- /@and@/@or@/@xor@ for integers, and @add@/@mul@/@minnumf@/@maxnumf@
-- /@minimumf@/@maximumf@ for floats. The default is @add@.
-- 
-- Example:
-- 
-- @
-- // Simple DOT product (K = 0).
-- \#contraction_accesses = [
--  affine_map\<(i) -> (i)>,
--  affine_map\<(i) -> (i)>,
--  affine_map\<(i) -> ()>
-- ]
-- \#contraction_trait = {
--   indexing_maps = \#contraction_accesses,
--   iterator_types = [\"reduction\"]
-- }
-- %3 = vector.contract \#contraction_trait %0, %1, %2
--   : vector\<10xf32>, vector\<10xf32> into f32
-- 
-- // 2D vector contraction with one contracting dimension (matmul, K = 2).
-- \#contraction_accesses = [
--   affine_map\<(i, j, k) -> (i, k)>,
--   affine_map\<(i, j, k) -> (k, j)>,
--   affine_map\<(i, j, k) -> (i, j)>
-- ]
-- \#contraction_trait = {
--   indexing_maps = \#contraction_accesses,
--   iterator_types = [\"parallel\", \"parallel\", \"reduction\"]
-- }
-- 
-- %3 = vector.contract \#contraction_trait %0, %1, %2
--   : vector\<4x3xf32>, vector\<3x7xf32> into vector\<4x7xf32>
-- 
-- // 4D to 3D vector contraction with two contracting dimensions and
-- // one batch dimension (K = 3).
-- \#contraction_accesses = [
--   affine_map\<(b0, f0, f1, c0, c1) -> (c0, b0, c1, f0)>,
--   affine_map\<(b0, f0, f1, c0, c1) -> (b0, c1, c0, f1)>,
--   affine_map\<(b0, f0, f1, c0, c1) -> (b0, f0, f1)>
-- ]
-- \#contraction_trait = {
--   indexing_maps = \#contraction_accesses,
--   iterator_types = [\"parallel\", \"parallel\", \"parallel\",
--                     \"reduction\", \"reduction\"]
-- }
-- 
-- %4 = vector.contract \#contraction_trait %0, %1, %2
--     : vector\<7x8x16x15xf32>, vector\<8x16x7x5xf32> into vector\<8x15x5xf32>
-- 
-- // Vector contraction with mixed typed. lhs/rhs have different element
-- // types than accumulator/result.
-- %5 = vector.contract \#contraction_trait %0, %1, %2
--   : vector\<10xf16>, vector\<10xf16> into f32
-- 
-- // Contract with max (K = 0).
-- \#contraction_accesses = [
--  affine_map\<(i) -> (i)>,
--  affine_map\<(i) -> (i)>,
--  affine_map\<(i) -> ()>
-- ]
-- \#contraction_trait = {
--   indexing_maps = \#contraction_accesses,
--   iterator_types = [\"reduction\"],
--   kind = \#vector.kind\<maxnumf>
-- }
-- %6 = vector.contract \#contraction_trait %0, %1, %2
--   : vector\<10xf32>, vector\<10xf32> into f32
-- @
--   

-- * create_mask
-- $create_mask
-- 
-- Creates and returns a vector mask where elements of the result vector
-- are set to \'0\' or \'1\', based on whether the element indices are contained
-- within a hyper-rectangular region specified by the operands. Specifically,
-- each operand specifies a range [0, operand-value) for a unique dimension in
-- the vector result. The conjunction of the operand ranges define a
-- hyper-rectangular region within which elements values are set to 1
-- (otherwise element values are set to 0). If operand-value is negative, it is
-- treated as if it were zero, and if it is greater than the corresponding
-- dimension size, it is treated as if it were equal to the dimension size.
-- 
-- Example:
-- 
-- @
-- // create a vector mask of size 4x3xi1 where elements in range
-- // 0 \<= row \<= 2 and 0 \<= col \<= 1 are set to 1 (others to 0).
-- %1 = vector.create_mask %c3, %c2 : vector\<4x3xi1>
-- 
-- print %1
--               columns
--             0    1    2
--           |------------
--         0 | 1    1    0
--   rows  1 | 1    1    0
--         2 | 1    1    0
--         3 | 0    0    0
-- @
--   

-- | A pattern for @vector.create_mask@.
pattern CreateMask :: () => () => Location -> Type -> [operand] -> AbstractOperation operand
pattern $bCreateMask :: Location -> Type -> [operand] -> AbstractOperation operand
$mCreateMask :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> [operand] -> r) -> (Void# -> r) -> r
CreateMask loc ty0 operands_  = Operation
          { opName = "vector.create_mask"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = operands_
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.create_mask@.
create_mask :: () => MonadBlockBuilder m => Type -> [Value] -> m Value
create_mask :: Type -> [Value] -> m Value
create_mask Type
ty0 [Value]
operands_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.create_mask"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([Value] -> [Name]
AST.operands [Value]
operands_)
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * deinterleave
-- $deinterleave
-- 
-- The deinterleave operation constructs two vectors from a single input
-- vector. The first result vector contains the elements from even indexes
-- of the input, and the second contains elements from odd indexes. This is
-- the inverse of a @vector.interleave@ operation.
-- 
-- Each output\'s trailing dimension is half of the size of the input
-- vector\'s trailing dimension. This operation requires the input vector
-- to have a rank > 0 and an even number of elements in its trailing
-- dimension.
-- 
-- The operation supports scalable vectors.
-- 
-- Example:
-- @
-- %0, %1 = vector.deinterleave %a
--            : vector\<8xi8> -> vector\<4xi8>
-- %2, %3 = vector.deinterleave %b
--            : vector\<2x8xi8> -> vector\<2x4xi8>
-- %4, %5 = vector.deinterleave %c
--            : vector\<2x8x4xi8> -> vector\<2x8x2xi8>
-- %6, %7 = vector.deinterleave %d
--            : vector\<[8]xf32> -> vector\<[4]xf32>
-- %8, %9 = vector.deinterleave %e
--            : vector\<2x[6]xf64> -> vector\<2x[3]xf64>
-- %10, %11 = vector.deinterleave %f
--            : vector\<2x4x[6]xf64> -> vector\<2x4x[3]xf64>
-- @
--       

-- | A pattern for @vector.deinterleave@.
pattern Deinterleave :: () => () => Location -> Type -> Type -> operand -> AbstractOperation operand
pattern $bDeinterleave :: Location -> Type -> Type -> operand -> AbstractOperation operand
$mDeinterleave :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> Type -> operand -> r) -> (Void# -> r) -> r
Deinterleave loc ty0 ty1 source_  = Operation
          { opName = "vector.deinterleave"
          , opLocation = loc
          , opResultTypes = Explicit [ty0, ty1]
          , opOperands = [source_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.deinterleave@.
deinterleave :: () => MonadBlockBuilder m => Type -> Type -> Value -> m [Value]
deinterleave :: Type -> Type -> Value -> m [Value]
deinterleave Type
ty0 Type
ty1 Value
source_   = do
  (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.deinterleave"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0, Type
ty1]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
source_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * expandload
-- $expandload
-- 
-- The expand load reads elements from memory into a vector as defined by a
-- base with indices and a mask vector. Expansion only applies to the innermost
-- dimension. When the mask is set, the next element is read from memory.
-- Otherwise, the corresponding element is taken from a pass-through vector.
-- Informally the semantics are:
-- 
-- @
-- index = i
-- result[0] := if mask[0] then base[index++] else pass_thru[0]
-- result[1] := if mask[1] then base[index++] else pass_thru[1]
-- etc.
-- @
-- 
-- Note that the index increment is done conditionally.
-- 
-- If a mask bit is set and the corresponding index is out-of-bounds for the
-- given base, the behavior is undefined. If a mask bit is not set, the value
-- comes from the pass-through vector regardless of the index, and the index is
-- allowed to be out-of-bounds.
-- 
-- The expand load can be used directly where applicable, or can be used
-- during progressively lowering to bring other memory operations closer to
-- hardware ISA support for an expand. The semantics of the operation closely
-- correspond to those of the @llvm.masked.expandload@
-- intrinsic.
-- 
-- Note, at the moment this Op is only available for fixed-width vectors.
-- 
-- Examples:
-- 
-- @
-- %0 = vector.expandload %base[%i], %mask, %pass_thru
--    : memref\<?xf32>, vector\<8xi1>, vector\<8xf32> into vector\<8xf32>
-- 
-- %1 = vector.expandload %base[%i, %j], %mask, %pass_thru
--    : memref\<?x?xf32>, vector\<16xi1>, vector\<16xf32> into vector\<16xf32>
-- @
--   

-- | A builder for @vector.expandload@.
expandload :: () => MonadBlockBuilder m => Type -> Value -> [Value] -> Value -> Value -> m Value
expandload :: Type -> Value -> [Value] -> Value -> Value -> m Value
expandload Type
ty0 Value
base_ [Value]
indices_ Value
mask_ Value
pass_thru_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.expandload"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_) [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
mask_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
pass_thru_)])
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * extractelement
-- $extractelement
-- 
-- Takes a 0-D or 1-D vector and a optional dynamic index position and
-- extracts the scalar at that position.
-- 
-- Note that this instruction resembles vector.extract, but is restricted to
-- 0-D and 1-D vectors and relaxed to dynamic indices.
-- If the vector is 0-D, the position must be std::nullopt.
-- 
-- 
-- It is meant to be closer to LLVM\'s version:
-- https://llvm.org/docs/LangRef.html\#extractelement-instruction
-- 
-- Example:
-- 
-- @
-- %c = arith.constant 15 : i32
-- %1 = vector.extractelement %0[%c : i32]: vector\<16xf32>
-- %2 = vector.extractelement %z[]: vector\<f32>
-- @
--   

-- | A builder for @vector.extractelement@.
extractelement :: () => MonadBlockBuilder m => Type -> Value -> Maybe Value -> m Value
extractelement :: Type -> Value -> Maybe Value -> m Value
extractelement Type
ty0 Value
vector_ Maybe Value
position_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.extractelement"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
vector_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ (Maybe Name -> [Name]
forall a. Maybe a -> [a]
Data.Maybe.maybeToList (Value -> Name
AST.operand (Value -> Name) -> Maybe Value -> Maybe Name
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> Maybe Value
position_)))
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * extract
-- $extract
-- 
-- Takes an n-D vector and a k-D position and extracts the (n-k)-D vector at
-- the proper position. Degenerates to an element type if n-k is zero.
-- 
-- Static and dynamic indices must be greater or equal to zero and less than
-- the size of the corresponding dimension. The result is undefined if any
-- index is out-of-bounds. The value @-1@ represents a poison index, which
-- specifies that the extracted element is poison.
-- 
-- Example:
-- 
-- @
-- %1 = vector.extract %0[3]: vector\<8x16xf32> from vector\<4x8x16xf32>
-- %2 = vector.extract %0[2, 1, 3]: f32 from vector\<4x8x16xf32>
-- %3 = vector.extract %1[]: vector\<f32> from vector\<f32>
-- %4 = vector.extract %0[%a, %b, %c]: f32 from vector\<4x8x16xf32>
-- %5 = vector.extract %0[2, %b]: vector\<16xf32> from vector\<4x8x16xf32>
-- %6 = vector.extract %10[-1, %c]: f32 from vector\<4x16xf32>
-- @
--   

-- * extract_strided_slice
-- $extract_strided_slice
-- 
-- Takes an n-D vector, k-D @offsets@ integer array attribute, a k-sized
-- @sizes@ integer array attribute, a k-sized @strides@ integer array
-- attribute and extracts the n-D subvector at the proper offset.
-- 
-- At the moment strides must contain only 1s.
-- 
-- Returns an n-D vector where the first k-D dimensions match the @sizes@
-- attribute. The returned subvector contains the elements starting at offset
-- @offsets@ and ending at @offsets + sizes@.
-- 
-- Example:
-- 
-- @
-- %1 = vector.extract_strided_slice %0
--     {offsets = [0, 2], sizes = [2, 4], strides = [1, 1]}:
--   vector\<4x8x16xf32> to vector\<2x4x16xf32>
-- 
-- // TODO: Evolve to a range form syntax similar to:
-- %1 = vector.extract_strided_slice %0[0:2:1][2:4:1]
--   vector\<4x8x16xf32> to vector\<2x4x16xf32>
-- @
-- 
-- TODO: Implement support for poison indices.
--   

pattern InternalExtractStridedSliceOpAttributes :: () => () => [Int] -> [Int] -> [Int] -> NamedAttributes
pattern $bInternalExtractStridedSliceOpAttributes :: [Int] -> [Int] -> [Int] -> Map Name Attribute
$mInternalExtractStridedSliceOpAttributes :: forall r.
Map Name Attribute
-> ([Int] -> [Int] -> [Int] -> r) -> (Void# -> r) -> r
InternalExtractStridedSliceOpAttributes offsets_ sizes_ strides_ <- ((\m -> (M.lookup "offsets" m, M.lookup "sizes" m, M.lookup "strides" m)) -> (Just (PatternUtil.I64ArrayAttr offsets_), Just (PatternUtil.I64ArrayAttr sizes_), Just (PatternUtil.I64ArrayAttr strides_)))
  where InternalExtractStridedSliceOpAttributes [Int]
offsets_ [Int]
sizes_ [Int]
strides_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ [(Name
"offsets", [Int] -> Attribute
PatternUtil.I64ArrayAttr [Int]
offsets_)] [(Name, Attribute)] -> [(Name, Attribute)] -> [(Name, Attribute)]
forall a. [a] -> [a] -> [a]
++ [(Name
"sizes", [Int] -> Attribute
PatternUtil.I64ArrayAttr [Int]
sizes_)] [(Name, Attribute)] -> [(Name, Attribute)] -> [(Name, Attribute)]
forall a. [a] -> [a] -> [a]
++ [(Name
"strides", [Int] -> Attribute
PatternUtil.I64ArrayAttr [Int]
strides_)]

-- | A pattern for @vector.extract_strided_slice@.
pattern ExtractStridedSlice :: () => () => Location -> Type -> operand -> [Int] -> [Int] -> [Int] -> AbstractOperation operand
pattern $bExtractStridedSlice :: Location
-> Type
-> operand
-> [Int]
-> [Int]
-> [Int]
-> AbstractOperation operand
$mExtractStridedSlice :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> [Int] -> [Int] -> [Int] -> r)
-> (Void# -> r)
-> r
ExtractStridedSlice loc ty0 vector_ offsets_ sizes_ strides_ = Operation
          { opName = "vector.extract_strided_slice"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [vector_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (InternalExtractStridedSliceOpAttributes offsets_ sizes_ strides_)
          }

-- | A builder for @vector.extract_strided_slice@.
extract_strided_slice :: () => MonadBlockBuilder m => Type -> Value -> [Int] -> [Int] -> [Int] -> m Value
extract_strided_slice :: Type -> Value -> [Int] -> [Int] -> [Int] -> m Value
extract_strided_slice Type
ty0 Value
vector_ [Int]
offsets_ [Int]
sizes_ [Int]
strides_  = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.extract_strided_slice"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
vector_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = ([Int] -> [Int] -> [Int] -> Map Name Attribute
InternalExtractStridedSliceOpAttributes [Int]
offsets_ [Int]
sizes_ [Int]
strides_)
          }))

-- * fma
-- $fma
-- 
-- Multiply-add expressions operate on n-D vectors and compute a fused
-- pointwise multiply-and-accumulate: @\$result = \$lhs * \$rhs + \$acc@.
-- All operands and result have the same vector type. The semantics
-- of the operation correspond to those of the @llvm.fma@
-- intrinsic. In the
-- particular case of lowering to LLVM, this is guaranteed to lower
-- to the @llvm.fma.*@ intrinsic.
-- 
-- Example:
-- 
-- @
-- %3 = vector.fma %0, %1, %2: vector\<8x16xf32>
-- @
--   

-- | A pattern for @vector.fma@.
pattern FMA :: () => () => Location -> Type -> operand -> operand -> operand -> AbstractOperation operand
pattern $bFMA :: Location
-> Type
-> operand
-> operand
-> operand
-> AbstractOperation operand
$mFMA :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> operand -> operand -> r)
-> (Void# -> r)
-> r
FMA loc ty0 lhs_ rhs_ acc_  = Operation
          { opName = "vector.fma"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [lhs_, rhs_, acc_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.fma@.
fma :: () => MonadBlockBuilder m => Type -> Value -> Value -> Value -> m Value
fma :: Type -> Value -> Value -> Value -> m Value
fma Type
ty0 Value
lhs_ Value
rhs_ Value
acc_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.fma"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
lhs_), (Value -> Name
AST.operand Value
rhs_), (Value -> Name
AST.operand Value
acc_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * flat_transpose
-- $flat_transpose
-- 
-- This is the counterpart of llvm.matrix.transpose in MLIR. It serves
-- the purposes of more progressive lowering and localized type conversion.
-- Higher levels typically lower matrix transpositions into \'vector.transpose\'
-- operations. Subsequent rewriting rule progressively lower these operations
-- into \'vector.flat_transpose\' operations to bring the operations closer
-- to the hardware ISA.
-- 
-- The @vector.flat_transpose@ op treats the 1-D input @matrix@ as
-- a 2-D matrix with \<rows> rows and \<columns> columns, and returns the
-- transposed matrix in flattened form in \'res\'.
-- 
-- Note, the corresponding LLVM intrinsic, @\@llvm.matrix.transpose.*@, does not
-- support scalable vectors. Hence, this Op is only available for fixed-width
-- vectors. Also see:
-- 
-- http://llvm.org/docs/LangRef.html\#llvm-matrix-transpose-intrinsic
-- 
-- Example:
-- 
-- @
-- %1 = vector.flat_transpose %0 {columns = 4 : i32, rows = 4 : i32}
--    : vector\<16xf32> -> vector\<16xf32>
-- @
--   

pattern InternalFlatTransposeOpAttributes :: () => () => Int -> Int -> NamedAttributes
pattern $bInternalFlatTransposeOpAttributes :: Int -> Int -> Map Name Attribute
$mInternalFlatTransposeOpAttributes :: forall r.
Map Name Attribute -> (Int -> Int -> r) -> (Void# -> r) -> r
InternalFlatTransposeOpAttributes rows_ columns_ <- ((\m -> (M.lookup "rows" m, M.lookup "columns" m)) -> (Just (IntegerAttr (IntegerType Signless 32) rows_), Just (IntegerAttr (IntegerType Signless 32) columns_)))
  where InternalFlatTransposeOpAttributes Int
rows_ Int
columns_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ [(Name
"rows", Type -> Int -> Attribute
IntegerAttr (Signedness -> UInt -> Type
IntegerType Signedness
Signless UInt
32) Int
rows_)] [(Name, Attribute)] -> [(Name, Attribute)] -> [(Name, Attribute)]
forall a. [a] -> [a] -> [a]
++ [(Name
"columns", Type -> Int -> Attribute
IntegerAttr (Signedness -> UInt -> Type
IntegerType Signedness
Signless UInt
32) Int
columns_)]

-- | A pattern for @vector.flat_transpose@.
pattern FlatTranspose :: () => () => Location -> Type -> operand -> Int -> Int -> AbstractOperation operand
pattern $bFlatTranspose :: Location
-> Type -> operand -> Int -> Int -> AbstractOperation operand
$mFlatTranspose :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> Int -> Int -> r)
-> (Void# -> r)
-> r
FlatTranspose loc ty0 matrix_ rows_ columns_ = Operation
          { opName = "vector.flat_transpose"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [matrix_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (InternalFlatTransposeOpAttributes rows_ columns_)
          }

-- | A builder for @vector.flat_transpose@.
flat_transpose :: () => MonadBlockBuilder m => Type -> Value -> Int -> Int -> m Value
flat_transpose :: Type -> Value -> Int -> Int -> m Value
flat_transpose Type
ty0 Value
matrix_ Int
rows_ Int
columns_  = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.flat_transpose"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
matrix_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Int -> Int -> Map Name Attribute
InternalFlatTransposeOpAttributes Int
rows_ Int
columns_)
          }))

-- * from_elements
-- $from_elements
-- 
-- This operation defines a vector from one or multiple scalar elements. The
-- number of elements must match the number of elements in the result type.
-- All elements must have the same type, which must match the element type of
-- the result vector type.
-- 
-- @elements@ are a flattened version of the result vector in row-major order.
-- 
-- Example:
-- 
-- @
-- // %f1
-- %0 = vector.from_elements %f1 : vector\<f32>
-- // [%f1, %f2]
-- %1 = vector.from_elements %f1, %f2 : vector\<2xf32>
-- // [[%f1, %f2, %f3], [%f4, %f5, %f6]]
-- %2 = vector.from_elements %f1, %f2, %f3, %f4, %f5, %f6 : vector\<2x3xf32>
-- // [[[%f1, %f2]], [[%f3, %f4]], [[%f5, %f6]]]
-- %3 = vector.from_elements %f1, %f2, %f3, %f4, %f5, %f6 : vector\<3x1x2xf32>
-- @
-- 
-- Note, scalable vectors are not supported.
--   

-- | A pattern for @vector.from_elements@.
pattern FromElements :: () => () => Location -> Type -> [operand] -> AbstractOperation operand
pattern $bFromElements :: Location -> Type -> [operand] -> AbstractOperation operand
$mFromElements :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> [operand] -> r) -> (Void# -> r) -> r
FromElements loc ty0 elements_  = Operation
          { opName = "vector.from_elements"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = elements_
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.from_elements@.
from_elements :: () => MonadBlockBuilder m => Type -> [Value] -> m Value
from_elements :: Type -> [Value] -> m Value
from_elements Type
ty0 [Value]
elements_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.from_elements"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([Value] -> [Name]
AST.operands [Value]
elements_)
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * gather
-- $gather
-- 
-- The gather operation returns an n-D vector whose elements are either loaded
-- from memory or ranked tensor, or taken from a pass-through vector, depending
-- on the values of an n-D mask vector.
-- If a mask bit is set, the corresponding result element is defined by the base
-- with indices and the n-D index vector (each index is a 1-D offset on the base).
-- Otherwise, the corresponding element is taken from the n-D pass-through vector.
-- Informally the semantics are:
-- @
-- result[0] := if mask[0] then base[index[0]] else pass_thru[0]
-- result[1] := if mask[1] then base[index[1]] else pass_thru[1]
-- etc.
-- @
-- 
-- If a mask bit is set and the corresponding index is out-of-bounds for the
-- given base, the behavior is undefined. If a mask bit is not set, the value
-- comes from the pass-through vector regardless of the index, and the index is
-- allowed to be out-of-bounds.
-- 
-- The gather operation can be used directly where applicable, or can be used
-- during progressively lowering to bring other memory operations closer to
-- hardware ISA support for a gather.
-- 
-- Examples:
-- 
-- @
-- %0 = vector.gather %base[%c0][%v], %mask, %pass_thru
--    : memref\<?xf32>, vector\<2x16xi32>, vector\<2x16xi1>, vector\<2x16xf32> into vector\<2x16xf32>
-- 
-- %1 = vector.gather %base[%i, %j][%v], %mask, %pass_thru
--    : memref\<16x16xf32>, vector\<16xi32>, vector\<16xi1>, vector\<16xf32> into vector\<16xf32>
-- @
--   

-- | A builder for @vector.gather@.
gather :: () => MonadBlockBuilder m => Type -> Value -> [Value] -> Value -> Value -> Value -> m Value
gather :: Type -> Value -> [Value] -> Value -> Value -> Value -> m Value
gather Type
ty0 Value
base_ [Value]
indices_ Value
index_vec_ Value
mask_ Value
pass_thru_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.gather"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_) [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
index_vec_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
mask_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
pass_thru_)])
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * insertelement
-- $insertelement
-- 
-- Takes a scalar source, a 0-D or 1-D destination vector and a dynamic index
-- position and inserts the source into the destination at the proper position.
-- 
-- Note that this instruction resembles vector.insert, but is restricted to 0-D
-- and 1-D vectors and relaxed to dynamic indices.
-- 
-- It is meant to be closer to LLVM\'s version:
-- https://llvm.org/docs/LangRef.html\#insertelement-instruction
-- 
-- Example:
-- 
-- @
-- %c = arith.constant 15 : i32
-- %f = arith.constant 0.0f : f32
-- %1 = vector.insertelement %f, %0[%c : i32]: vector\<16xf32>
-- %2 = vector.insertelement %f, %z[]: vector\<f32>
-- @
--   

-- | A builder for @vector.insertelement@.
insertelement :: () => MonadBlockBuilder m => Type -> Value -> Value -> Maybe Value -> m Value
insertelement :: Type -> Value -> Value -> Maybe Value -> m Value
insertelement Type
ty0 Value
source_ Value
dest_ Maybe Value
position_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.insertelement"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
source_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
dest_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ (Maybe Name -> [Name]
forall a. Maybe a -> [a]
Data.Maybe.maybeToList (Value -> Name
AST.operand (Value -> Name) -> Maybe Value -> Maybe Name
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> Maybe Value
position_)))
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * insert
-- $insert
-- 
-- Takes an n-D source vector, an (n+k)-D destination vector and a k-D position
-- and inserts the n-D source into the (n+k)-D destination at the proper
-- position. Degenerates to a scalar or a 0-d vector source type when n = 0.
-- 
-- Static and dynamic indices must be greater or equal to zero and less than
-- the size of the corresponding dimension. The result is undefined if any
-- index is out-of-bounds. The value @-1@ represents a poison index, which
-- specifies that the resulting vector is poison.
-- 
-- Example:
-- 
-- @
-- %2 = vector.insert %0, %1[3] : vector\<8x16xf32> into vector\<4x8x16xf32>
-- %5 = vector.insert %3, %4[2, 1, 3] : f32 into vector\<4x8x16xf32>
-- %8 = vector.insert %6, %7[] : f32 into vector\<f32>
-- %11 = vector.insert %9, %10[%a, %b, %c] : vector\<f32> into vector\<4x8x16xf32>
-- %12 = vector.insert %4, %10[2, %b] : vector\<16xf32> into vector\<4x8x16xf32>
-- %13 = vector.insert %20, %1[-1, %c] : f32 into vector\<4x16xf32>
-- @
--   

-- * insert_strided_slice
-- $insert_strided_slice
-- 
-- Takes a k-D source vector, an n-D destination vector (n >= k), n-sized
-- @offsets@ integer array attribute, a k-sized @strides@ integer array attribute
-- and inserts the k-D source vector as a strided subvector at the proper offset
-- into the n-D destination vector.
-- 
-- At the moment strides must contain only 1s.
-- 
-- Returns an n-D vector that is a copy of the n-D destination vector in which
-- the last k-D dimensions contain the k-D source vector elements strided at
-- the proper location as specified by the offsets.
-- 
-- Example:
-- 
-- @
-- %2 = vector.insert_strided_slice %0, %1
--     {offsets = [0, 0, 2], strides = [1, 1]}:
--   vector\<2x4xf32> into vector\<16x4x8xf32>
-- @
--   

pattern InternalInsertStridedSliceOpAttributes :: () => () => [Int] -> [Int] -> NamedAttributes
pattern $bInternalInsertStridedSliceOpAttributes :: [Int] -> [Int] -> Map Name Attribute
$mInternalInsertStridedSliceOpAttributes :: forall r.
Map Name Attribute -> ([Int] -> [Int] -> r) -> (Void# -> r) -> r
InternalInsertStridedSliceOpAttributes offsets_ strides_ <- ((\m -> (M.lookup "offsets" m, M.lookup "strides" m)) -> (Just (PatternUtil.I64ArrayAttr offsets_), Just (PatternUtil.I64ArrayAttr strides_)))
  where InternalInsertStridedSliceOpAttributes [Int]
offsets_ [Int]
strides_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ [(Name
"offsets", [Int] -> Attribute
PatternUtil.I64ArrayAttr [Int]
offsets_)] [(Name, Attribute)] -> [(Name, Attribute)] -> [(Name, Attribute)]
forall a. [a] -> [a] -> [a]
++ [(Name
"strides", [Int] -> Attribute
PatternUtil.I64ArrayAttr [Int]
strides_)]

-- | A pattern for @vector.insert_strided_slice@.
pattern InsertStridedSlice :: () => () => Location -> Type -> operand -> operand -> [Int] -> [Int] -> AbstractOperation operand
pattern $bInsertStridedSlice :: Location
-> Type
-> operand
-> operand
-> [Int]
-> [Int]
-> AbstractOperation operand
$mInsertStridedSlice :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> operand -> [Int] -> [Int] -> r)
-> (Void# -> r)
-> r
InsertStridedSlice loc ty0 source_ dest_ offsets_ strides_ = Operation
          { opName = "vector.insert_strided_slice"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [source_, dest_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (InternalInsertStridedSliceOpAttributes offsets_ strides_)
          }

-- | A builder for @vector.insert_strided_slice@.
insert_strided_slice :: () => MonadBlockBuilder m => Type -> Value -> Value -> [Int] -> [Int] -> m Value
insert_strided_slice :: Type -> Value -> Value -> [Int] -> [Int] -> m Value
insert_strided_slice Type
ty0 Value
source_ Value
dest_ [Int]
offsets_ [Int]
strides_  = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.insert_strided_slice"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
source_), (Value -> Name
AST.operand Value
dest_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = ([Int] -> [Int] -> Map Name Attribute
InternalInsertStridedSliceOpAttributes [Int]
offsets_ [Int]
strides_)
          }))

-- * interleave
-- $interleave
-- 
-- The interleave operation constructs a new vector by interleaving the
-- elements from the trailing (or final) dimension of two input vectors,
-- returning a new vector where the trailing dimension is twice the size.
-- 
-- Note that for the n-D case this differs from the interleaving possible with
-- @vector.shuffle@, which would only operate on the leading dimension.
-- 
-- Another key difference is this operation supports scalable vectors, though
-- currently a general LLVM lowering is limited to the case where only the
-- trailing dimension is scalable.
-- 
-- Example:
-- @
-- %a = arith.constant dense\<[0, 1]> : vector\<2xi32>
-- %b = arith.constant dense\<[2, 3]> : vector\<2xi32>
-- // The value of @%0@ is @[0, 2, 1, 3]@.
-- %0 = vector.interleave %a, %b : vector\<2xi32> -> vector\<4xi32>
-- 
-- // Examples showing allowed input and result types.
-- %1 = vector.interleave %c, %d : vector\<f16> -> vector\<2xf16>
-- %2 = vector.interleave %e, %f : vector\<6x3xf32> -> vector\<6x6xf32>
-- %3 = vector.interleave %g, %h : vector\<[4]xi32> -> vector\<[8]xi32>
-- %4 = vector.interleave %i, %j : vector\<2x4x[2]xf64> -> vector\<2x4x[4]xf64>
-- @
--   

-- | A pattern for @vector.interleave@.
pattern Interleave :: () => () => Location -> Type -> operand -> operand -> AbstractOperation operand
pattern $bInterleave :: Location -> Type -> operand -> operand -> AbstractOperation operand
$mInterleave :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> operand -> r)
-> (Void# -> r)
-> r
Interleave loc ty0 lhs_ rhs_  = Operation
          { opName = "vector.interleave"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [lhs_, rhs_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.interleave@.
interleave :: () => MonadBlockBuilder m => Type -> Value -> Value -> m Value
interleave :: Type -> Value -> Value -> m Value
interleave Type
ty0 Value
lhs_ Value
rhs_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.interleave"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
lhs_), (Value -> Name
AST.operand Value
rhs_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * load
-- $load
-- 
-- The \'vector.load\' operation reads an n-D slice of memory into an n-D
-- vector. It takes a \'base\' memref, an index for each memref dimension and a
-- result vector type as arguments. It returns a value of the result vector
-- type. The \'base\' memref and indices determine the start memory address from
-- which to read. Each index provides an offset for each memref dimension
-- based on the element type of the memref. The shape of the result vector
-- type determines the shape of the slice read from the start memory address.
-- The elements along each dimension of the slice are strided by the memref
-- strides. When loading more than 1 element, only unit strides are allowed
-- along the most minor memref dimension. These constraints guarantee that
-- elements read along the first dimension of the slice are contiguous in
-- memory.
-- 
-- The memref element type can be a scalar or a vector type. If the memref
-- element type is a scalar, it should match the element type of the result
-- vector. If the memref element type is vector, it should match the result
-- vector type.
-- 
-- Example: 0-D vector load on a scalar memref.
-- @
-- %result = vector.load %base[%i, %j] : memref\<100x100xf32>, vector\<f32>
-- @
-- 
-- Example: 1-D vector load on a scalar memref.
-- @
-- %result = vector.load %base[%i, %j] : memref\<100x100xf32>, vector\<8xf32>
-- @
-- 
-- Example: 1-D vector load on a vector memref.
-- @
-- %result = vector.load %memref[%i, %j] : memref\<200x100xvector\<8xf32>>, vector\<8xf32>
-- @
-- 
-- Example:  2-D vector load on a scalar memref.
-- @
-- %result = vector.load %memref[%i, %j] : memref\<200x100xf32>, vector\<4x8xf32>
-- @
-- 
-- Example:  2-D vector load on a vector memref.
-- @
-- %result = vector.load %memref[%i, %j] : memref\<200x100xvector\<4x8xf32>>, vector\<4x8xf32>
-- @
-- 
-- Representation-wise, the \'vector.load\' operation permits out-of-bounds
-- reads. Support and implementation of out-of-bounds vector loads is
-- target-specific. No assumptions should be made on the value of elements
-- loaded out of bounds. Not all targets may support out-of-bounds vector
-- loads.
-- 
-- Example:  Potential out-of-bound vector load.
-- @
-- %result = vector.load %memref[%index] : memref\<?xf32>, vector\<8xf32>
-- @
-- 
-- Example:  Explicit out-of-bound vector load.
-- @
-- %result = vector.load %memref[%c0] : memref\<7xf32>, vector\<8xf32>
-- @
--   

pattern OptionalBoolAttr :: Maybe Bool -> Maybe Attribute
pattern $bOptionalBoolAttr :: Maybe Bool -> Maybe Attribute
$mOptionalBoolAttr :: forall r. Maybe Attribute -> (Maybe Bool -> r) -> (Void# -> r) -> r
OptionalBoolAttr x <- ((\case Just (BoolAttr y) -> Just y; Nothing -> Nothing) -> x)
  where OptionalBoolAttr Maybe Bool
x = case Maybe Bool
x of Just Bool
y -> Attribute -> Maybe Attribute
forall a. a -> Maybe a
Just (Bool -> Attribute
BoolAttr Bool
y); Maybe Bool
Nothing -> Maybe Attribute
forall a. Maybe a
Nothing

pattern InternalLoadOpAttributes :: () => () => Maybe Bool -> NamedAttributes
pattern $bInternalLoadOpAttributes :: Maybe Bool -> Map Name Attribute
$mInternalLoadOpAttributes :: forall r.
Map Name Attribute -> (Maybe Bool -> r) -> (Void# -> r) -> r
InternalLoadOpAttributes nontemporal_ <- ((\m -> (M.lookup "nontemporal" m)) -> (OptionalBoolAttr nontemporal_))
  where InternalLoadOpAttributes Maybe Bool
nontemporal_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ (Maybe (Name, Attribute) -> [(Name, Attribute)]
forall a. Maybe a -> [a]
Data.Maybe.maybeToList (Maybe (Name, Attribute) -> [(Name, Attribute)])
-> Maybe (Name, Attribute) -> [(Name, Attribute)]
forall a b. (a -> b) -> a -> b
$ (Name
"nontemporal",) (Attribute -> (Name, Attribute))
-> Maybe Attribute -> Maybe (Name, Attribute)
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> Maybe Bool -> Maybe Attribute
OptionalBoolAttr Maybe Bool
nontemporal_)

-- | A builder for @vector.load@.
load :: () => MonadBlockBuilder m => Type -> Value -> [Value] -> Maybe Bool -> m Value
load :: Type -> Value -> [Value] -> Maybe Bool -> m Value
load Type
ty0 Value
base_ [Value]
indices_ Maybe Bool
nontemporal_  = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.load"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_))
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Maybe Bool -> Map Name Attribute
InternalLoadOpAttributes Maybe Bool
nontemporal_)
          }))

-- * mask
-- $mask
-- 
-- The @vector.mask@ is a @MaskingOpInterface@ operation that predicates the
-- execution of another operation. It takes an @i1@ vector mask and an
-- optional passthru vector as arguments.
-- 
-- A implicitly @vector.yield@-terminated region encloses the operation to be
-- masked. Values used within the region are captured from above. Only one
-- *maskable* operation can be masked with a @vector.mask@ operation at a time.
-- An operation is *maskable* if it implements the @MaskableOpInterface@. The
-- terminator yields all results of the maskable operation to the result of
-- this operation.
-- 
-- The vector mask argument holds a bit for each vector lane and determines
-- which vector lanes should execute the maskable operation and which ones
-- should not. The @vector.mask@ operation returns the value produced by the
-- masked execution of the nested operation, if any. The masked-off lanes in
-- the result vector are taken from the corresponding lanes of the pass-thru
-- argument, if provided, or left unmodified, otherwise. At this point, 0-D
-- vectors are not supported by @vector.mask@. They may be supported in the
-- future.
-- 
-- The @vector.mask@ operation does not prescribe how a maskable operation
-- should be masked or how a masked operation should be lowered. Masking
-- constraints and some semantic details are provided by each maskable
-- operation through the @MaskableOpInterface@. Lowering of masked operations
-- is implementation defined. For instance, scalarizing the masked operation
-- or executing the operation for the masked-off lanes are valid lowerings as
-- long as the execution of masked-off lanes does not change the observable
-- behavior of the program.
-- 
-- Examples:
-- 
-- @
--   %0 = vector.mask %mask { vector.reduction \<add>, %a : vector\<8xi32> into i32 } : vector\<8xi1> -> i32
-- @
-- 
-- @
--   %0 = vector.mask %mask, %passthru { arith.divsi %a, %b : vector\<8xi32> } : vector\<8xi1> -> vector\<8xi32>
-- @
-- 
-- @
--   vector.mask %mask { vector.transfer_write %val, %t0[%idx] : vector\<16xf32>, memref\<?xf32> } : vector\<16xi1>
-- @
-- 
-- @
--   vector.mask %mask { vector.transfer_write %val, %t0[%idx] : vector\<16xf32>, tensor\<?xf32> } : vector\<16xi1> -> tensor\<?xf32>
-- @
--   

-- | A builder for @vector.mask@.
mask :: () => MonadBlockBuilder m => [Type] -> Value -> Maybe Value -> RegionBuilderT m () -> m Value
mask :: [Type] -> Value -> Maybe Value -> RegionBuilderT m () -> m Value
mask [Type]
ty0 Value
mask_ Maybe Value
passthru_  RegionBuilderT m ()
maskRegion_Builder = do
  Region
maskRegion_ <- RegionBuilderT m () -> m Region
forall (m :: * -> *). Monad m => RegionBuilderT m () -> m Region
AST.buildRegion RegionBuilderT m ()
maskRegion_Builder
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.mask"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit ([Type]
ty0)
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
mask_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ (Maybe Name -> [Name]
forall a. Maybe a -> [a]
Data.Maybe.maybeToList (Value -> Name
AST.operand (Value -> Name) -> Maybe Value -> Maybe Name
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> Maybe Value
passthru_)))
          , opRegions :: [Region]
opRegions = [Region
maskRegion_]
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * maskedload
-- $maskedload
-- 
-- The masked load reads elements from memory into a vector as defined
-- by a base with indices and a mask vector. When the mask is set, the
-- element is read from memory. Otherwise, the corresponding element is taken
-- from a pass-through vector. Informally the semantics are:
-- @
-- result[0] := if mask[0] then base[i + 0] else pass_thru[0]
-- result[1] := if mask[1] then base[i + 1] else pass_thru[1]
-- etc.
-- @
-- 
-- If a mask bit is set and the corresponding index is out-of-bounds for the
-- given base, the behavior is undefined. If a mask bit is not set, the value
-- comes from the pass-through vector regardless of the index, and the index is
-- allowed to be out-of-bounds.
-- 
-- The masked load can be used directly where applicable, or can be used
-- during progressively lowering to bring other memory operations closer to
-- hardware ISA support for a masked load. The semantics of the operation
-- closely correspond to those of the @llvm.masked.load@
-- intrinsic.
-- 
-- Examples:
-- 
-- @
-- %0 = vector.maskedload %base[%i], %mask, %pass_thru
--    : memref\<?xf32>, vector\<8xi1>, vector\<8xf32> into vector\<8xf32>
-- 
-- %1 = vector.maskedload %base[%i, %j], %mask, %pass_thru
--    : memref\<?x?xf32>, vector\<16xi1>, vector\<16xf32> into vector\<16xf32>
-- @
--   

-- | A builder for @vector.maskedload@.
maskedload :: () => MonadBlockBuilder m => Type -> Value -> [Value] -> Value -> Value -> m Value
maskedload :: Type -> Value -> [Value] -> Value -> Value -> m Value
maskedload Type
ty0 Value
base_ [Value]
indices_ Value
mask_ Value
pass_thru_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.maskedload"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_) [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
mask_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
pass_thru_)])
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * maskedstore
-- $maskedstore
-- 
-- The masked store operation writes elements from a vector into memory
-- as defined by a base with indices and a mask vector. When the mask is
-- set, the corresponding element from the vector is written to memory. Otherwise,
-- no action is taken for the element. Informally the semantics are:
-- @
-- if (mask[0]) base[i+0] = value[0]
-- if (mask[1]) base[i+1] = value[1]
-- etc.
-- @
-- 
-- If a mask bit is set and the corresponding index is out-of-bounds for the
-- given base, the behavior is undefined. If a mask bit is not set, no value
-- is stored regardless of the index, and the index is allowed to be
-- out-of-bounds.
-- 
-- The masked store can be used directly where applicable, or can be used
-- during progressively lowering to bring other memory operations closer to
-- hardware ISA support for a masked store. The semantics of the operation
-- closely correspond to those of the @llvm.masked.store@
-- intrinsic.
-- 
-- Examples:
-- 
-- @
-- vector.maskedstore %base[%i], %mask, %value
--   : memref\<?xf32>, vector\<8xi1>, vector\<8xf32>
-- 
-- vector.maskedstore %base[%i, %j], %mask, %value
--   : memref\<?x?xf32>, vector\<16xi1>, vector\<16xf32>
-- @
--   

-- | A builder for @vector.maskedstore@.
maskedstore :: () => MonadBlockBuilder m => Value -> [Value] -> Value -> Value -> m ()
maskedstore :: Value -> [Value] -> Value -> Value -> m ()
maskedstore  Value
base_ [Value]
indices_ Value
mask_ Value
valueToStore_   = do
  m [Value] -> m ()
forall (f :: * -> *) a. Functor f => f a -> f ()
Control.Monad.void (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.maskedstore"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit []
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_) [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
mask_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
valueToStore_)])
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * matrix_multiply
-- $matrix_multiply
-- 
-- This is the counterpart of llvm.matrix.multiply in MLIR. It serves the
-- purposes of more progressive lowering and localized type conversion.
-- Higher levels typically lower matrix multiplications into \'vector.contract\'
-- operations. Subsequent rewriting rule progressively lower these operations
-- into \'vector.matrix_multiply\' operations to bring the operations closer
-- to the hardware ISA.
-- 
-- The ‘vector.matrix_multiply’ op treats @lhs@ as matrix with \<lhs_rows> rows
-- and \<lhs_columns> columns, @rhs@ as matrix with \<lhs_columns> rows and
-- \<rhs_columns> and multiplies them. The result matrix is returned embedded in
-- the result vector.
-- 
-- Note, the corresponding LLVM intrinsic, @\@llvm.matrix.multiply.*@, does not
-- support scalable vectors. Hence, this Op is only available for fixed-width
-- vectors. Also see:
-- 
-- http://llvm.org/docs/LangRef.html\#llvm-matrix-multiply-intrinsic
-- 
-- Example:
-- 
-- @
-- %C = vector.matrix_multiply %A, %B
--   { lhs_rows = 4: i32, lhs_columns = 16: i32 , rhs_columns = 3: i32 } :
--   (vector\<64xf64>, vector\<48xf64>) -> vector\<12xf64>
-- @
--   

pattern InternalMatmulOpAttributes :: () => () => Int -> Int -> Int -> NamedAttributes
pattern $bInternalMatmulOpAttributes :: Int -> Int -> Int -> Map Name Attribute
$mInternalMatmulOpAttributes :: forall r.
Map Name Attribute -> (Int -> Int -> Int -> r) -> (Void# -> r) -> r
InternalMatmulOpAttributes lhs_rows_ lhs_columns_ rhs_columns_ <- ((\m -> (M.lookup "lhs_rows" m, M.lookup "lhs_columns" m, M.lookup "rhs_columns" m)) -> (Just (IntegerAttr (IntegerType Signless 32) lhs_rows_), Just (IntegerAttr (IntegerType Signless 32) lhs_columns_), Just (IntegerAttr (IntegerType Signless 32) rhs_columns_)))
  where InternalMatmulOpAttributes Int
lhs_rows_ Int
lhs_columns_ Int
rhs_columns_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ [(Name
"lhs_rows", Type -> Int -> Attribute
IntegerAttr (Signedness -> UInt -> Type
IntegerType Signedness
Signless UInt
32) Int
lhs_rows_)] [(Name, Attribute)] -> [(Name, Attribute)] -> [(Name, Attribute)]
forall a. [a] -> [a] -> [a]
++ [(Name
"lhs_columns", Type -> Int -> Attribute
IntegerAttr (Signedness -> UInt -> Type
IntegerType Signedness
Signless UInt
32) Int
lhs_columns_)] [(Name, Attribute)] -> [(Name, Attribute)] -> [(Name, Attribute)]
forall a. [a] -> [a] -> [a]
++ [(Name
"rhs_columns", Type -> Int -> Attribute
IntegerAttr (Signedness -> UInt -> Type
IntegerType Signedness
Signless UInt
32) Int
rhs_columns_)]

-- | A pattern for @vector.matrix_multiply@.
pattern Matmul :: () => () => Location -> Type -> operand -> operand -> Int -> Int -> Int -> AbstractOperation operand
pattern $bMatmul :: Location
-> Type
-> operand
-> operand
-> Int
-> Int
-> Int
-> AbstractOperation operand
$mMatmul :: forall r operand.
AbstractOperation operand
-> (Location
    -> Type -> operand -> operand -> Int -> Int -> Int -> r)
-> (Void# -> r)
-> r
Matmul loc ty0 lhs_ rhs_ lhs_rows_ lhs_columns_ rhs_columns_ = Operation
          { opName = "vector.matrix_multiply"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [lhs_, rhs_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (InternalMatmulOpAttributes lhs_rows_ lhs_columns_ rhs_columns_)
          }

-- | A builder for @vector.matrix_multiply@.
matrix_multiply :: () => MonadBlockBuilder m => Type -> Value -> Value -> Int -> Int -> Int -> m Value
matrix_multiply :: Type -> Value -> Value -> Int -> Int -> Int -> m Value
matrix_multiply Type
ty0 Value
lhs_ Value
rhs_ Int
lhs_rows_ Int
lhs_columns_ Int
rhs_columns_  = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.matrix_multiply"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
lhs_), (Value -> Name
AST.operand Value
rhs_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Int -> Int -> Int -> Map Name Attribute
InternalMatmulOpAttributes Int
lhs_rows_ Int
lhs_columns_ Int
rhs_columns_)
          }))

-- * multi_reduction
-- $multi_reduction
-- 
-- Reduces an n-D vector into an (n-k)-D vector (or a scalar when k == n)
-- using the given operation: @add@/@mul@/@minsi@/@minui@/@maxsi@/@maxui@
-- /@and@/@or@/@xor@ for integers, and @add@/@mul@/@minnumf@/@maxnumf@/@minimumf@
-- /@maximumf@ for floats.
-- Takes an initial accumulator operand.
-- 
-- Example:
-- 
-- @
-- %1 = vector.multi_reduction \<add>, %0, %acc0 [1, 3] :
--   vector\<4x8x16x32xf32> to vector\<4x16xf32>
-- %2 = vector.multi_reduction \<add>, %1, %acc1 [0, 1] :
--   vector\<4x16xf32> to f32
-- @
--   

-- * outerproduct
-- $outerproduct
-- 
-- Takes 2 1-D vectors and returns the 2-D vector containing the outer-product,
-- as illustrated below:
-- @
--  outer |   [c, d]
--  ------+------------
--    [a, | [ [a*c, a*d],
--     b] |   [b*c, b*d] ]
-- @
-- This operation also accepts a 1-D vector lhs and a scalar rhs. In this
-- case a simple AXPY operation is performed, which returns a 1-D vector.
-- @
--     [a, b] * c = [a*c, b*c]
-- @
-- 
-- An optional extra vector argument with the same shape as the output
-- vector may be specified in which case the operation returns the sum of
-- the outer-product and the extra vector. In this multiply-accumulate
-- scenario for floating-point arguments, the rounding mode is enforced
-- by guaranteeing that a fused-multiply add operation is emitted. When
-- lowered to the LLVMIR dialect, this form emits @llvm.intr.fma@, which
-- is guaranteed to lower to actual @fma@ instructions on x86.
-- 
-- An optional kind attribute may be specified to be: @add@/@mul@/@minsi@
-- /@minui@/@maxsi@/@maxui@/@and@/@or@/@xor@ for integers, and @add@/@mul@
-- /@minnumf@/@maxnumf@/@minimumf@/@maximumf@ for floats. The default is
-- @add@.
-- 
-- Example:
-- 
-- @
-- %2 = vector.outerproduct %0, %1: vector\<4xf32>, vector\<8xf32>
-- return %2: vector\<4x8xf32>
-- 
-- %3 = vector.outerproduct %0, %1, %2:
--   vector\<4xf32>, vector\<8xf32>, vector\<4x8xf32>
-- return %3: vector\<4x8xf32>
-- 
-- %4 = vector.outerproduct %0, %1, %2 {kind = \#vector.kind\<maxnumf>}:
--   vector\<4xf32>, vector\<8xf32>, vector\<4x8xf32>
-- return %3: vector\<4x8xf32>
-- 
-- %6 = vector.outerproduct %4, %5: vector\<10xf32>, f32
-- return %6: vector\<10xf32>
-- 
-- @
--   

-- | A builder for @vector.outerproduct@.
outerproduct :: () => MonadBlockBuilder m => Type -> Value -> Value -> Maybe Value -> m Value
outerproduct :: Type -> Value -> Value -> Maybe Value -> m Value
outerproduct Type
ty0 Value
lhs_ Value
rhs_ Maybe Value
acc_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.outerproduct"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
lhs_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
rhs_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ (Maybe Name -> [Name]
forall a. Maybe a -> [a]
Data.Maybe.maybeToList (Value -> Name
AST.operand (Value -> Name) -> Maybe Value -> Maybe Name
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> Maybe Value
acc_)))
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * print
-- $print
-- 
-- Prints the source vector (or scalar) to stdout in a human-readable format
-- (for testing and debugging). No return value.
-- 
-- Example:
-- 
-- @
-- %v = arith.constant dense\<0.0> : vector\<4xf32>
-- vector.print %v : vector\<4xf32>
-- @
-- 
-- When lowered to LLVM, the vector print is decomposed into elementary
-- printing method calls that at runtime will yield:
-- 
-- @
-- ( 0.0, 0.0, 0.0, 0.0 )
-- @
-- 
-- This is printed to stdout via a small runtime support library, which only
-- needs to provide a few printing methods (single value for all data
-- types, opening/closing bracket, comma, newline).
-- 
-- By default @vector.print@ adds a newline after the vector, but this can be
-- controlled by the @punctuation@ attribute. For example, to print a comma
-- after instead do:
-- 
-- @
-- vector.print %v : vector\<4xf32> punctuation \<comma>
-- @
-- 
-- Note that it is possible to use the punctuation attribute alone. The
-- following will print a single newline:
-- 
-- @
-- vector.print punctuation \<newline>
-- @
-- 
-- Additionally, to aid with debugging and testing @vector.print@ can also
-- print constant strings:
-- 
-- @
-- vector.print str \"Hello, World!\"
-- @
--   

-- | A builder for @vector.print@.
print :: () => MonadBlockBuilder m => Maybe Value -> m ()
print :: Maybe Value -> m ()
print  Maybe Value
source_   = do
  m [Value] -> m ()
forall (f :: * -> *) a. Functor f => f a -> f ()
Control.Monad.void (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.print"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit []
          , opOperands :: [Name]
opOperands = ((Maybe Name -> [Name]
forall a. Maybe a -> [a]
Data.Maybe.maybeToList (Value -> Name
AST.operand (Value -> Name) -> Maybe Value -> Maybe Name
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> Maybe Value
source_)))
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * reduction
-- $reduction
-- 
-- Reduces an 1-D vector \"horizontally\" into a scalar using the given
-- operation: @add@/@mul@/@minsi@/@minui@/@maxsi@/@maxui@/@and@/@or@/@xor@ for
-- integers, and @add@/@mul@/@minnumf@/@maxnumf@/@minimumf@/@maximumf@ for
-- floats. Reductions also allow an optional fused accumulator.
-- 
-- Note that these operations are restricted to 1-D vectors to remain
-- close to the corresponding LLVM intrinsics:
-- 
-- http://llvm.org/docs/LangRef.html\#vector-reduction-intrinsics
-- 
-- Example:
-- 
-- @
-- %1 = vector.reduction \<add>, %0 : vector\<16xf32> into f32
-- 
-- %3 = vector.reduction \<xor>, %2 : vector\<4xi32> into i32
-- 
-- %4 = vector.reduction \<mul>, %0, %1 : vector\<16xf32> into f32
-- @
--   

-- * scalable.extract
-- $scalable.extract
-- 
-- Takes rank-1 source vector and a position @pos@ within the source
-- vector, and extracts a subvector starting from that position.
-- 
-- The extraction position must be a multiple of the minimum size of the result
-- vector. For the operation to be well defined, the destination vector must
-- fit within the source vector from the specified position. Since the source
-- vector is scalable and its runtime length is unknown, the validity of the
-- operation can\'t be verified nor guaranteed at compile time.
-- 
-- Example:
-- 
-- @
-- %1 = vector.scalable.extract %0[8] : vector\<4xf32> from vector\<[8]xf32>
-- %3 = vector.scalable.extract %2[0] : vector\<[4]xf32> from vector\<[8]xf32>
-- @
-- 
-- Invalid example:
-- @
-- %1 = vector.scalable.extract %0[5] : vector\<4xf32> from vector\<[16]xf32>
-- @
--   

pattern InternalScalableExtractOpAttributes :: () => () => Int -> NamedAttributes
pattern $bInternalScalableExtractOpAttributes :: Int -> Map Name Attribute
$mInternalScalableExtractOpAttributes :: forall r. Map Name Attribute -> (Int -> r) -> (Void# -> r) -> r
InternalScalableExtractOpAttributes pos_ <- ((\m -> (M.lookup "pos" m)) -> (Just (IntegerAttr (IntegerType Signless 64) pos_)))
  where InternalScalableExtractOpAttributes Int
pos_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ [(Name
"pos", Type -> Int -> Attribute
IntegerAttr (Signedness -> UInt -> Type
IntegerType Signedness
Signless UInt
64) Int
pos_)]

-- | A pattern for @vector.scalable.extract@.
pattern ScalableExtract :: () => () => Location -> Type -> operand -> Int -> AbstractOperation operand
pattern $bScalableExtract :: Location -> Type -> operand -> Int -> AbstractOperation operand
$mScalableExtract :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> Int -> r) -> (Void# -> r) -> r
ScalableExtract loc ty0 source_ pos_ = Operation
          { opName = "vector.scalable.extract"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [source_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (InternalScalableExtractOpAttributes pos_)
          }

-- | A builder for @vector.scalable.extract@.
scalable_extract :: () => MonadBlockBuilder m => Type -> Value -> Int -> m Value
scalable_extract :: Type -> Value -> Int -> m Value
scalable_extract Type
ty0 Value
source_ Int
pos_  = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.scalable.extract"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
source_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Int -> Map Name Attribute
InternalScalableExtractOpAttributes Int
pos_)
          }))

-- * scalable.insert
-- $scalable.insert
-- 
-- This operations takes a rank-1 fixed-length or scalable subvector and
-- inserts it within the destination scalable vector starting from the
-- position specificed by @pos@. If the source vector is scalable, the
-- insertion position will be scaled by the runtime scaling factor of the
-- source subvector.
-- 
-- The insertion position must be a multiple of the minimum size of the source
-- vector. For the operation to be well defined, the source vector must fit in
-- the destination vector from the specified position. Since the destination
-- vector is scalable and its runtime length is unknown, the validity of the
-- operation can\'t be verified nor guaranteed at compile time.
-- 
-- Example:
-- 
-- @
-- %2 = vector.scalable.insert %0, %1[8] : vector\<4xf32> into vector\<[16]xf32>
-- %5 = vector.scalable.insert %3, %4[0] : vector\<8xf32> into vector\<[4]xf32>
-- %8 = vector.scalable.insert %6, %7[0] : vector\<[4]xf32> into vector\<[8]xf32>
-- @
-- 
-- Invalid example:
-- @
-- %2 = vector.scalable.insert %0, %1[5] : vector\<4xf32> into vector\<[16]xf32>
-- @
--   

pattern InternalScalableInsertOpAttributes :: () => () => Int -> NamedAttributes
pattern $bInternalScalableInsertOpAttributes :: Int -> Map Name Attribute
$mInternalScalableInsertOpAttributes :: forall r. Map Name Attribute -> (Int -> r) -> (Void# -> r) -> r
InternalScalableInsertOpAttributes pos_ <- ((\m -> (M.lookup "pos" m)) -> (Just (IntegerAttr (IntegerType Signless 64) pos_)))
  where InternalScalableInsertOpAttributes Int
pos_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ [(Name
"pos", Type -> Int -> Attribute
IntegerAttr (Signedness -> UInt -> Type
IntegerType Signedness
Signless UInt
64) Int
pos_)]

-- | A pattern for @vector.scalable.insert@.
pattern ScalableInsert :: () => () => Location -> Type -> operand -> operand -> Int -> AbstractOperation operand
pattern $bScalableInsert :: Location
-> Type -> operand -> operand -> Int -> AbstractOperation operand
$mScalableInsert :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> operand -> Int -> r)
-> (Void# -> r)
-> r
ScalableInsert loc ty0 source_ dest_ pos_ = Operation
          { opName = "vector.scalable.insert"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [source_, dest_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (InternalScalableInsertOpAttributes pos_)
          }

-- | A builder for @vector.scalable.insert@.
scalable_insert :: () => MonadBlockBuilder m => Type -> Value -> Value -> Int -> m Value
scalable_insert :: Type -> Value -> Value -> Int -> m Value
scalable_insert Type
ty0 Value
source_ Value
dest_ Int
pos_  = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.scalable.insert"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
source_), (Value -> Name
AST.operand Value
dest_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Int -> Map Name Attribute
InternalScalableInsertOpAttributes Int
pos_)
          }))

-- * scan
-- $scan
-- 
-- Performs an inclusive/exclusive scan on an n-D vector along a single
-- dimension returning an n-D result vector using the given
-- operation (@add@/@mul@/@minsi@/@minui@/@maxsi@/@maxui@/@and@/@or@/@xor@ for
-- integers, and @add@/@mul@/@minnumf@/@maxnumf@/@minimumf@/@maximumf@ for
-- floats), and a specified value for the initial value. The operator returns
-- the result of scan as well as the result of the last reduction in the scan.
-- 
-- Example:
-- 
-- @
-- %1:2 = vector.scan \<add>, %0, %acc {inclusive = false, reduction_dim = 1 : i64} :
--   vector\<4x8x16x32xf32>, vector\<4x16x32xf32>
-- @
--   

-- * scatter
-- $scatter
-- 
-- The scatter operation stores elements from a 1-D vector into memory as
-- defined by a base with indices and an additional 1-D index vector, but
-- only if the corresponding bit in a 1-D mask vector is set. Otherwise, no
-- action is taken for that element. Informally the semantics are:
-- @
-- if (mask[0]) base[index[0]] = value[0]
-- if (mask[1]) base[index[1]] = value[1]
-- etc.
-- @
-- 
-- If a mask bit is set and the corresponding index is out-of-bounds for the
-- given base, the behavior is undefined. If a mask bit is not set, no value
-- is stored regardless of the index, and the index is allowed to be
-- out-of-bounds.
-- 
-- If the index vector contains two or more duplicate indices, the behavior is
-- undefined. Underlying implementation may enforce strict sequential
-- semantics.
-- TODO: always enforce strict sequential semantics?
-- 
-- The scatter operation can be used directly where applicable, or can be used
-- during progressively lowering to bring other memory operations closer to
-- hardware ISA support for a scatter. The semantics of the operation closely
-- correspond to those of the @llvm.masked.scatter@
-- intrinsic.
-- 
-- Examples:
-- 
-- @
-- vector.scatter %base[%c0][%v], %mask, %value
--     : memref\<?xf32>, vector\<16xi32>, vector\<16xi1>, vector\<16xf32>
-- 
-- vector.scatter %base[%i, %j][%v], %mask, %value
--     : memref\<16x16xf32>, vector\<16xi32>, vector\<16xi1>, vector\<16xf32>
-- @
--   

-- | A builder for @vector.scatter@.
scatter :: () => MonadBlockBuilder m => Value -> [Value] -> Value -> Value -> Value -> m ()
scatter :: Value -> [Value] -> Value -> Value -> Value -> m ()
scatter  Value
base_ [Value]
indices_ Value
index_vec_ Value
mask_ Value
valueToStore_   = do
  m [Value] -> m ()
forall (f :: * -> *) a. Functor f => f a -> f ()
Control.Monad.void (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.scatter"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit []
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_) [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
index_vec_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
mask_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
valueToStore_)])
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * shape_cast
-- $shape_cast
-- 
-- The shape_cast operation casts between an n-D source vector shape and
-- a k-D result vector shape (the element type remains the same).
-- 
-- If reducing rank (n > k), result dimension sizes must be a product
-- of contiguous source dimension sizes.
-- If expanding rank (n \< k), source dimensions must factor into a
-- contiguous sequence of destination dimension sizes.
-- Each source dim is expanded (or contiguous sequence of source dims combined)
-- in source dimension list order (i.e. 0 \<= i \< n), to produce a contiguous
-- sequence of result dims (or a single result dim), in result dimension list
-- order (i.e. 0 \<= j \< k). The product of all source dimension sizes and all
-- result dimension sizes must match.
-- 
-- It is currently assumed that this operation does not require moving data,
-- and that it will be folded away before lowering vector operations.
-- 
-- There is an exception to the folding expectation when targeting
-- llvm.intr.matrix operations. We need a type conversion back and forth from a
-- 2-D MLIR vector to a 1-D flattened LLVM vector.shape_cast lowering to LLVM
-- is supported in that particular case, for now.
-- 
-- Example:
-- 
-- @
-- // Example casting to a lower vector rank.
-- %1 = vector.shape_cast %0 : vector\<5x1x4x3xf32> to vector\<20x3xf32>
-- 
-- // Example casting to a higher vector rank.
-- %3 = vector.shape_cast %2 : vector\<10x12x8xf32> to vector\<5x2x3x4x8xf32>
-- 
-- @
--   

-- | A pattern for @vector.shape_cast@.
pattern ShapeCast :: () => () => Location -> Type -> operand -> AbstractOperation operand
pattern $bShapeCast :: Location -> Type -> operand -> AbstractOperation operand
$mShapeCast :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> r) -> (Void# -> r) -> r
ShapeCast loc ty0 source_  = Operation
          { opName = "vector.shape_cast"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [source_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.shape_cast@.
shape_cast :: () => MonadBlockBuilder m => Type -> Value -> m Value
shape_cast :: Type -> Value -> m Value
shape_cast Type
ty0 Value
source_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.shape_cast"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
source_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * shuffle
-- $shuffle
-- 
-- The shuffle operation constructs a permutation (or duplication) of elements
-- from two input vectors, returning a vector with the same element type as
-- the input and a length that is the same as the shuffle mask. The two input
-- vectors must have the same element type, same rank, and trailing dimension
-- sizes and shuffles their values in the leading dimension (which may differ
-- in size) according to the given mask. The legality rules are:
-- * the two operands must have the same element type as the result
--   - Either, the two operands and the result must have the same
--     rank and trailing dimension sizes, viz. given two k-D operands
--             v1 : \<s_1 x s_2 x .. x s_k x type> and
--             v2 : \<t_1 x t_2 x .. x t_k x type>
--     we have s_i = t_i for all 1 \< i \<= k
--   - Or, the two operands must be 0-D vectors and the result is a 1-D vector.
-- * the mask length equals the leading dimension size of the result
-- * numbering the input vector indices left to right across the operands, all
--   mask values must be within range, viz. given two k-D operands v1 and v2
--   above, all mask values are in the range [0,s_1+t_1). The value @-1@
--   represents a poison mask value, which specifies that the selected element
--   is poison.
-- 
-- Note, scalable vectors are not supported.
-- 
-- Example:
-- 
-- @
-- %0 = vector.shuffle %a, %b[0, 3]
--            : vector\<2xf32>, vector\<2xf32>       ; yields vector\<2xf32>
-- %1 = vector.shuffle %c, %b[0, 1, 2]
--            : vector\<2x16xf32>, vector\<1x16xf32> ; yields vector\<3x16xf32>
-- %2 = vector.shuffle %a, %b[3, 2, 1, 0]
--            : vector\<2xf32>, vector\<2xf32>       ; yields vector\<4xf32>
-- %3 = vector.shuffle %a, %b[0, 1]
--            : vector\<f32>, vector\<f32>           ; yields vector\<2xf32>
-- %4 = vector.shuffle %a, %b[0, 4, -1, -1, -1, -1]
--            : vector\<4xf32>, vector\<4xf32>       ; yields vector\<6xf32>
-- @
--   

-- * splat
-- $splat
-- 
-- Broadcast the operand to all elements of the result vector. The operand is
-- required to be of integer/index/float type.
-- 
-- Example:
-- 
-- @
-- %s = arith.constant 10.1 : f32
-- %t = vector.splat %s : vector\<8x16xf32>
-- @
--   

-- | A pattern for @vector.splat@.
pattern Splat :: () => () => Location -> Type -> operand -> AbstractOperation operand
pattern $bSplat :: Location -> Type -> operand -> AbstractOperation operand
$mSplat :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> r) -> (Void# -> r) -> r
Splat loc ty0 input_  = Operation
          { opName = "vector.splat"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [input_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.splat@.
splat :: () => MonadBlockBuilder m => Type -> Value -> m Value
splat :: Type -> Value -> m Value
splat Type
ty0 Value
input_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.splat"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
input_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * step
-- $step
-- 
-- A @step@ operation produces an index vector, i.e. a 1-D vector of values of
-- index type that represents a linear sequence from 0 to N-1, where N is the
-- number of elements in the @result@ vector.
-- 
-- Supports fixed-width and scalable vectors.
-- 
-- Examples:
-- 
-- @
-- %0 = vector.step : vector\<4xindex> ; [0, 1, 2, 3]
-- %1 = vector.step : vector\<[4]xindex> ; [0, 1, .., \<vscale * 4 - 1>]
-- @
--   

-- | A pattern for @vector.step@.
pattern Step :: () => () => Location -> Type -> AbstractOperation operand
pattern $bStep :: Location -> Type -> AbstractOperation operand
$mStep :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> r) -> (Void# -> r) -> r
Step loc ty0   = Operation
          { opName = "vector.step"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = []
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.step@.
step :: () => MonadBlockBuilder m => Type -> m Value
step :: Type -> m Value
step Type
ty0    = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.step"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = []
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * store
-- $store
-- 
-- The \'vector.store\' operation writes an n-D vector to an n-D slice of memory.
-- It takes the vector value to be stored, a \'base\' memref and an index for
-- each memref dimension. The \'base\' memref and indices determine the start
-- memory address from which to write. Each index provides an offset for each
-- memref dimension based on the element type of the memref. The shape of the
-- vector value to store determines the shape of the slice written from the
-- start memory address. The elements along each dimension of the slice are
-- strided by the memref strides. When storing more than 1 element, only unit
-- strides are allowed along the most minor memref dimension. These constraints
-- guarantee that elements written along the first dimension of the slice are
-- contiguous in memory.
-- 
-- The memref element type can be a scalar or a vector type. If the memref
-- element type is a scalar, it should match the element type of the value
-- to store. If the memref element type is vector, it should match the type
-- of the value to store.
-- 
-- Example: 0-D vector store on a scalar memref.
-- @
-- vector.store %valueToStore, %memref[%i, %j] : memref\<200x100xf32>, vector\<f32>
-- @
-- 
-- Example: 1-D vector store on a scalar memref.
-- @
-- vector.store %valueToStore, %memref[%i, %j] : memref\<200x100xf32>, vector\<8xf32>
-- @
-- 
-- Example: 1-D vector store on a vector memref.
-- @
-- vector.store %valueToStore, %memref[%i, %j] : memref\<200x100xvector\<8xf32>>, vector\<8xf32>
-- @
-- 
-- Example:  2-D vector store on a scalar memref.
-- @
-- vector.store %valueToStore, %memref[%i, %j] : memref\<200x100xf32>, vector\<4x8xf32>
-- @
-- 
-- Example:  2-D vector store on a vector memref.
-- @
-- vector.store %valueToStore, %memref[%i, %j] : memref\<200x100xvector\<4x8xf32>>, vector\<4x8xf32>
-- @
-- 
-- Representation-wise, the \'vector.store\' operation permits out-of-bounds
-- writes. Support and implementation of out-of-bounds vector stores are
-- target-specific. No assumptions should be made on the memory written out of
-- bounds. Not all targets may support out-of-bounds vector stores.
-- 
-- Example:  Potential out-of-bounds vector store.
-- @
-- vector.store %valueToStore, %memref[%index] : memref\<?xf32>, vector\<8xf32>
-- @
-- 
-- Example:  Explicit out-of-bounds vector store.
-- @
-- vector.store %valueToStore, %memref[%c0] : memref\<7xf32>, vector\<8xf32>
-- @
--   

pattern InternalStoreOpAttributes :: () => () => Maybe Bool -> NamedAttributes
pattern $bInternalStoreOpAttributes :: Maybe Bool -> Map Name Attribute
$mInternalStoreOpAttributes :: forall r.
Map Name Attribute -> (Maybe Bool -> r) -> (Void# -> r) -> r
InternalStoreOpAttributes nontemporal_ <- ((\m -> (M.lookup "nontemporal" m)) -> (OptionalBoolAttr nontemporal_))
  where InternalStoreOpAttributes Maybe Bool
nontemporal_ = [(Name, Attribute)] -> Map Name Attribute
forall k a. Ord k => [(k, a)] -> Map k a
M.fromList ([(Name, Attribute)] -> Map Name Attribute)
-> [(Name, Attribute)] -> Map Name Attribute
forall a b. (a -> b) -> a -> b
$ (Maybe (Name, Attribute) -> [(Name, Attribute)]
forall a. Maybe a -> [a]
Data.Maybe.maybeToList (Maybe (Name, Attribute) -> [(Name, Attribute)])
-> Maybe (Name, Attribute) -> [(Name, Attribute)]
forall a b. (a -> b) -> a -> b
$ (Name
"nontemporal",) (Attribute -> (Name, Attribute))
-> Maybe Attribute -> Maybe (Name, Attribute)
forall (f :: * -> *) a b. Functor f => (a -> b) -> f a -> f b
<$> Maybe Bool -> Maybe Attribute
OptionalBoolAttr Maybe Bool
nontemporal_)

-- | A builder for @vector.store@.
store :: () => MonadBlockBuilder m => Value -> Value -> [Value] -> Maybe Bool -> m ()
store :: Value -> Value -> [Value] -> Maybe Bool -> m ()
store  Value
valueToStore_ Value
base_ [Value]
indices_ Maybe Bool
nontemporal_  = do
  m [Value] -> m ()
forall (f :: * -> *) a. Functor f => f a -> f ()
Control.Monad.void (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.store"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit []
          , opOperands :: [Name]
opOperands = ([(Value -> Name
AST.operand Value
valueToStore_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ [(Value -> Name
AST.operand Value
base_)] [Name] -> [Name] -> [Name]
forall a. [a] -> [a] -> [a]
++ ([Value] -> [Name]
AST.operands [Value]
indices_))
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Maybe Bool -> Map Name Attribute
InternalStoreOpAttributes Maybe Bool
nontemporal_)
          }))

-- * transfer_read
-- $transfer_read
-- 
-- The @vector.transfer_read@ op performs a read from a slice within a
-- MemRef or a Ranked
-- Tensor supplied as its first operand
-- into a vector of the same base elemental type.
-- 
-- A memref/tensor operand with vector element type, must have its vector
-- element type match a suffix (shape and element type) of the vector (e.g.
-- memref\<3x2x6x4x3xf32>, vector\<1x1x4x3xf32>).
-- 
-- The slice is further defined by a full-rank index within the MemRef/Tensor,
-- supplied as the operands @[1 .. 1 + rank(memref/tensor))@ that defines the
-- starting point of the transfer (e.g. @%A[%i0, %i1, %i2]@).
-- 
-- The permutation_map attribute is an
-- affine-map which specifies the transposition on the
-- slice to match the vector shape. The permutation map may be implicit and
-- omitted from parsing and printing if it is the canonical minor identity map
-- (i.e. if it does not permute or broadcast any dimension).
-- 
-- The size of the slice is specified by the size of the vector, given as the
-- return type.
-- 
-- An SSA value @padding@ of the same elemental type as the MemRef/Tensor is
-- provided to specify a fallback value in the case of out-of-bounds accesses
-- and/or masking.
-- 
-- An optional SSA value @mask@ may be specified to mask out elements read from
-- the MemRef/Tensor. The @mask@ type is an @i1@ vector with a shape that
-- matches how elements are read from the MemRef/Tensor, *before* any
-- permutation or broadcasting. Elements whose corresponding mask element is
-- @0@ are masked out and replaced with @padding@.
-- 
-- For every vector dimension, the boolean array attribute @in_bounds@
-- specifies if the transfer is guaranteed to be within the source bounds. If
-- set to \"false\", accesses (including the starting point) may run
-- out-of-bounds along the respective vector dimension as the index increases.
-- Non-vector dimensions *must* always be in-bounds. The @in_bounds@ array
-- length has to be equal to the vector rank. This attribute has a default
-- value: @false@ (i.e. \"out-of-bounds\"). When skipped in the textual IR, the
-- default value is assumed. Similarly, the OP printer will omit this
-- attribute when all dimensions are out-of-bounds (i.e. the default value is
-- used).
-- 
-- A @vector.transfer_read@ can be lowered to a simple load if all dimensions
-- are specified to be within bounds and no @mask@ was specified.
-- 
-- This operation is called \'read\' by opposition to \'load\' because the
-- super-vector granularity is generally not representable with a single
-- hardware register. A @vector.transfer_read@ is thus a mid-level abstraction
-- that supports super-vectorization with non-effecting padding for full-tile
-- only operations.
-- 
-- More precisely, let\'s dive deeper into the permutation_map for the following
-- MLIR:
-- 
-- @
-- vector.transfer_read %A[%expr1, %expr2, %expr3, %expr4]
--   { permutation_map : (d0,d1,d2,d3) -> (d2,0,d0) } :
--   memref\<?x?x?x?xf32>, vector\<3x4x5xf32>
-- @
-- 
-- This operation always reads a slice starting at @%A[%expr1, %expr2, %expr3,
-- %expr4]@. The size of the slice can be inferred from the resulting vector
-- shape and walking back through the permutation map: 3 along d2 and 5 along
-- d0, so the slice is: @%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@.
-- 
-- A notional lowering of vector.transfer_read could generate code resembling:
-- 
-- @
-- // %expr1, %expr2, %expr3, %expr4 defined before this point
-- // alloc a temporary buffer for performing the \"gather\" of the slice.
-- %tmp = memref.alloc() : memref\<vector\<3x4x5xf32>>
-- for %i = 0 to 3 {
--   affine.for %j = 0 to 4 {
--     affine.for %k = 0 to 5 {
--       // Note that this load does not involve %j.
--       %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] : memref\<?x?x?x?xf32>
--       // Update the temporary gathered slice with the individual element
--       %slice = memref.load %tmp : memref\<vector\<3x4x5xf32>> -> vector\<3x4x5xf32>
--       %updated = vector.insert %a, %slice[%i, %j, %k] : f32 into vector\<3x4x5xf32>
--       memref.store %updated, %tmp : memref\<vector\<3x4x5xf32>>
-- }}}
-- // At this point we gathered the elements from the original
-- // memref into the desired vector layout, stored in the @%tmp@ allocation.
-- %vec = memref.load %tmp : memref\<vector\<3x4x5xf32>> -> vector\<3x4x5xf32>
-- @
-- 
-- On a GPU one could then map @i@, @j@, @k@ to blocks and threads. Notice that
-- the temporary storage footprint could conceptually be only @3 * 5@ values but
-- @3 * 4 * 5@ values are actually transferred between @%A@ and @%tmp@.
-- 
-- Alternatively, if a notional vector broadcast operation were available, we
-- could avoid the loop on @%j@ and the lowered code would resemble:
-- 
-- @
-- // %expr1, %expr2, %expr3, %expr4 defined before this point
-- %tmp = memref.alloc() : memref\<vector\<3x4x5xf32>>
-- for %i = 0 to 3 {
--   affine.for %k = 0 to 5 {
--     %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] : memref\<?x?x?x?xf32>
--     %slice = memref.load %tmp : memref\<vector\<3x4x5xf32>> -> vector\<3x4x5xf32>
--     // Here we only store to the first element in dimension one
--     %updated = vector.insert %a, %slice[%i, 0, %k] : f32 into vector\<3x4x5xf32>
--     memref.store %updated, %tmp : memref\<vector\<3x4x5xf32>>
-- }}
-- // At this point we gathered the elements from the original
-- // memref into the desired vector layout, stored in the @%tmp@ allocation.
-- // However we haven\'t replicated them alongside the first dimension, we need
-- // to broadcast now.
-- %partialVec = load %tmp : memref\<vector\<3x4x5xf32>> -> vector\<3x4x5xf32>
-- %vec = broadcast %tmpvec, 1 : vector\<3x4x5xf32>
-- @
-- 
-- where @broadcast@ broadcasts from element 0 to all others along the
-- specified dimension. This time, the number of loaded element is @3 * 5@
-- values.
-- An additional @1@ broadcast is required. On a GPU this broadcast could be
-- implemented using a warp-shuffle if loop @j@ were mapped to @threadIdx.x@.
-- 
-- Syntax
-- @
-- operation ::= ssa-id @=@ @vector.transfer_read@ ssa-use-list
--   @{@ attribute-entry @} :@ memref-type @,@ vector-type
-- @
-- 
-- Example:
-- 
-- @
-- // Read the slice @%A[%i0, %i1:%i1+256, %i2:%i2+32]@ into vector\<32x256xf32>
-- // and pad with %f0 to handle the boundary case:
-- %f0 = arith.constant 0.0f : f32
-- affine.for %i0 = 0 to %0 {
--   affine.for %i1 = 0 to %1 step 256 {
--     affine.for %i2 = 0 to %2 step 32 {
--       %v = vector.transfer_read %A[%i0, %i1, %i2], (%f0)
--            {permutation_map: (d0, d1, d2) -> (d2, d1)} :
--            memref\<?x?x?xf32>, vector\<32x256xf32>
-- }}}
-- 
-- // or equivalently (rewrite with vector.transpose)
-- %f0 = arith.constant 0.0f : f32
-- affine.for %i0 = 0 to %0 {
--   affine.for %i1 = 0 to %1 step 256 {
--     affine.for %i2 = 0 to %2 step 32 {
--       %v0 = vector.transfer_read %A[%i0, %i1, %i2], (%f0)
--            {permutation_map: (d0, d1, d2) -> (d1, d2)} :
--            memref\<?x?x?xf32>, vector\<256x32xf32>
--       %v = vector.transpose %v0, [1, 0] :
--           vector\<256x32xf32> to vector\<32x256f32>
-- }}}
-- 
-- // Read the slice @%A[%i0, %i1]@ (i.e. the element @%A[%i0, %i1]@) into
-- // vector\<128xf32>. The underlying implementation will require a 1-D vector
-- // broadcast:
-- affine.for %i0 = 0 to %0 {
--   affine.for %i1 = 0 to %1 {
--     %3 = vector.transfer_read %A[%i0, %i1]
--          {permutation_map: (d0, d1) -> (0)} :
--          memref\<?x?xf32>, vector\<128xf32>
--   }
-- }
-- 
-- // Read from a memref with vector element type.
-- %4 = vector.transfer_read %arg1[%c3, %c3], %vf0
--   {permutation_map = (d0, d1)->(d0, d1)}
--     : memref\<?x?xvector\<4x3xf32>>, vector\<1x1x4x3xf32>
-- 
-- // Read from a tensor with vector element type.
-- %4 = vector.transfer_read %arg1[%c3, %c3], %vf0
--   {permutation_map = (d0, d1)->(d0, d1)}
--     : tensor\<?x?xvector\<4x3xf32>>, vector\<1x1x4x3xf32>
-- 
-- // Special encoding for 0-d transfer with 0-d tensor/memref, vector shape
-- // {1} and permutation_map () -> (0).
-- %0 = vector.transfer_read %arg0[], %f0 {permutation_map = affine_map\<()->(0)>} :
--   tensor\<f32>, vector\<1xf32>
-- @
--   

-- * transfer_write
-- $transfer_write
-- 
-- The @vector.transfer_write@ op performs a write from a
-- vector, supplied as its first operand, into a
-- slice within a MemRef or a Ranked
-- Tensor of the same base elemental type,
-- supplied as its second 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.
-- 
-- The slice is further defined by a full-rank index within the MemRef/Tensor,
-- supplied as the operands @[2 .. 2 + rank(memref/tensor))@ that defines the
-- starting point of the transfer (e.g. @%A[%i0, %i1, %i2, %i3]@).
-- 
-- The permutation_map attribute is an
-- affine-map which specifies the transposition on the
-- slice to match the vector shape. The permutation map may be implicit and
-- omitted from parsing and printing if it is the canonical minor identity map
-- (i.e. if it does not permute any dimension). In contrast to @transfer_read@,
-- write ops cannot have broadcast dimensions.
-- 
-- The size of the slice is specified by the size of the vector.
-- 
-- An optional SSA value @mask@ may be specified to mask out elements written
-- to the MemRef/Tensor. The @mask@ type is an @i1@ vector with a shape that
-- matches how elements are written into the MemRef/Tensor, *after* applying
-- any permutation. Elements whose corresponding mask element is @0@ are
-- masked out.
-- 
-- For every vector dimension, the boolean array attribute @in_bounds@
-- specifies if the transfer is guaranteed to be within the source bounds. If
-- set to \"false\", accesses (including the starting point) may run
-- out-of-bounds along the respective vector dimension as the index increases.
-- Non-vector dimensions *must* always be in-bounds. The @in_bounds@ array
-- length has to be equal to the vector rank. This attribute has a default
-- value: @false@ (i.e. \"out-of-bounds\"). When skipped in the textual IR, the
-- default value is assumed. Similarly, the OP printer will omit this
-- attribute when all dimensions are out-of-bounds (i.e. the default value is
-- used).
-- 
--  A @vector.transfer_write@ can be lowered to a simple store if all
--  dimensions are specified to be within bounds and no @mask@ was specified.
-- 
-- This operation is called \'write\' by opposition to \'store\' because the
-- super-vector granularity is generally not representable with a single
-- hardware register. A @vector.transfer_write@ is thus a
-- mid-level abstraction that supports super-vectorization with non-effecting
-- padding for full-tile-only code. It is the responsibility of
-- @vector.transfer_write@\'s implementation to ensure the memory writes are
-- valid. Different lowerings may be pertinent depending on the hardware
-- support.
-- 
-- Example:
-- 
-- @
-- // write vector\<16x32x64xf32> into the slice
-- //   @%A[%i0, %i1:%i1+32, %i2:%i2+64, %i3:%i3+16]@:
-- for %i0 = 0 to %0 {
--   affine.for %i1 = 0 to %1 step 32 {
--     affine.for %i2 = 0 to %2 step 64 {
--       affine.for %i3 = 0 to %3 step 16 {
--         %val = @ssa-value@ : vector\<16x32x64xf32>
--         vector.transfer_write %val, %A[%i0, %i1, %i2, %i3]
--           {permutation_map: (d0, d1, d2, d3) -> (d3, d1, d2)} :
--           vector\<16x32x64xf32>, memref\<?x?x?x?xf32>
-- }}}}
-- 
-- // or equivalently (rewrite with vector.transpose)
-- for %i0 = 0 to %0 {
--   affine.for %i1 = 0 to %1 step 32 {
--     affine.for %i2 = 0 to %2 step 64 {
--       affine.for %i3 = 0 to %3 step 16 {
--         %val = @ssa-value@ : vector\<16x32x64xf32>
--         %valt = vector.transpose %val, [1, 2, 0] :
--               vector\<16x32x64xf32> -> vector\<32x64x16xf32>
--         vector.transfer_write %valt, %A[%i0, %i1, %i2, %i3]
--           {permutation_map: (d0, d1, d2, d3) -> (d1, d2, d3)} :
--           vector\<32x64x16xf32>, memref\<?x?x?x?xf32>
-- }}}}
-- 
-- // write to a memref with vector element type.
-- vector.transfer_write %4, %arg1[%c3, %c3]
--   {permutation_map = (d0, d1)->(d0, d1)}
--     : vector\<1x1x4x3xf32>, memref\<?x?xvector\<4x3xf32>>
-- 
-- // return a tensor where the vector is inserted into the source tensor.
-- %5 = vector.transfer_write %4, %arg1[%c3, %c3]
--   {permutation_map = (d0, d1)->(d0, d1)}
--     : vector\<1x1x4x3xf32>, tensor\<?x?xvector\<4x3xf32>>
-- 
-- // Special encoding for 0-d transfer with 0-d tensor/memref, vector shape
-- // {1} and permutation_map () -> (0).
-- %1 = vector.transfer_write %0, %arg0[] {permutation_map = affine_map\<()->(0)>} :
--   vector\<1xf32>, tensor\<f32>
-- @
--   

-- * transpose
-- $transpose
-- 
-- Takes a n-D vector and returns the transposed n-D vector defined by
-- the permutation of ranks in the n-sized integer array attribute (in case
-- of 0-D vectors the array attribute must be empty).
-- 
-- In the operation
-- 
-- @
-- %1 = vector.transpose %0, [i_1, .., i_n]
--   : vector\<d_1 x .. x d_n x f32>
--   to vector\<d_trans[0] x .. x d_trans[n-1] x f32>
-- @
-- 
-- the @permutation@ array [i_1, .., i_n] must be a permutation of [0, .., n-1].
-- 
-- Example:
-- 
-- @
-- %1 = vector.transpose %0, [1, 0] : vector\<2x3xf32> to vector\<3x2xf32>
-- 
--  [ [a, b, c],       [ [a, d],
--    [d, e, f] ]  ->    [b, e],
--                       [c, f] ]
-- @
--   

-- * type_cast
-- $type_cast
-- 
-- Performs a conversion from a memref with scalar element to a memref with a
-- *single* vector element, copying the shape of the memref to the vector. This
-- is the minimal viable operation that is required to makeke
-- super-vectorization operational. It can be seen as a special case of the
-- @view@ operation but scoped in the super-vectorization context.
-- 
-- Example:
-- 
-- @
-- %A  = memref.alloc() : memref\<5x4x3xf32>
-- %VA = vector.type_cast %A : memref\<5x4x3xf32> to memref\<vector\<5x4x3xf32>>
-- @
--   

-- | A pattern for @vector.type_cast@.
pattern TypeCast :: () => () => Location -> Type -> operand -> AbstractOperation operand
pattern $bTypeCast :: Location -> Type -> operand -> AbstractOperation operand
$mTypeCast :: forall r operand.
AbstractOperation operand
-> (Location -> Type -> operand -> r) -> (Void# -> r) -> r
TypeCast loc ty0 memref_  = Operation
          { opName = "vector.type_cast"
          , opLocation = loc
          , opResultTypes = Explicit [ty0]
          , opOperands = [memref_]
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.type_cast@.
type_cast :: () => MonadBlockBuilder m => Type -> Value -> m Value
type_cast :: Type -> Value -> m Value
type_cast Type
ty0 Value
memref_   = do
  ([Value] -> Value) -> m [Value] -> m Value
forall (m :: * -> *) a1 r. Monad m => (a1 -> r) -> m a1 -> m r
Control.Monad.liftM [Value] -> Value
forall a. [a] -> a
Prelude.head (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.type_cast"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit [Type
ty0]
          , opOperands :: [Name]
opOperands = [(Value -> Name
AST.operand Value
memref_)]
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))

-- * yield
-- $yield
-- 
-- \"vector.yield\" yields an SSA value from the Vector dialect op region and
-- terminates the regions. The semantics of how the values are yielded is
-- defined by the parent operation.
-- If \"vector.yield\" has any operands, the operands must correspond to the
-- parent operation\'s results.
-- If the parent operation defines no value the vector.yield may be omitted
-- when printing the region.
--   

-- | A pattern for @vector.yield@.
pattern Yield :: () => () => Location -> [operand] -> AbstractOperation operand
pattern $bYield :: Location -> [operand] -> AbstractOperation operand
$mYield :: forall r operand.
AbstractOperation operand
-> (Location -> [operand] -> r) -> (Void# -> r) -> r
Yield loc  operands_  = Operation
          { opName = "vector.yield"
          , opLocation = loc
          , opResultTypes = Explicit []
          , opOperands = operands_
          , opRegions = []
          , opSuccessors = []
          , opAttributes = (NoAttrs)
          }

-- | A builder for @vector.yield@.
yield :: () => MonadBlockBuilder m => [Value] -> m EndOfBlock
yield :: [Value] -> m EndOfBlock
yield  [Value]
operands_   = do
  m [Value] -> m ()
forall (f :: * -> *) a. Functor f => f a -> f ()
Control.Monad.void (Operation -> m [Value]
forall (m :: * -> *). MonadBlockBuilder m => Operation -> m [Value]
AST.emitOp (Operation :: forall operand.
Name
-> Location
-> ResultTypes
-> [operand]
-> [Region]
-> [Name]
-> Map Name Attribute
-> AbstractOperation operand
Operation
          { opName :: Name
opName = Name
"vector.yield"
          , opLocation :: Location
opLocation = Location
UnknownLocation
          , opResultTypes :: ResultTypes
opResultTypes = [Type] -> ResultTypes
Explicit []
          , opOperands :: [Name]
opOperands = ([Value] -> [Name]
AST.operands [Value]
operands_)
          , opRegions :: [Region]
opRegions = []
          , opSuccessors :: [Name]
opSuccessors = []
          , opAttributes :: Map Name Attribute
opAttributes = (Map Name Attribute
NoAttrs)
          }))
  m EndOfBlock
forall (m :: * -> *). Monad m => m EndOfBlock
AST.terminateBlock