#include "mlir/Dialect/NVGPU/TransformOps/NVGPUTransformOps.h"
#include "mlir/Analysis/SliceAnalysis.h"
#include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
#include "mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Arith/Utils/Utils.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialectDecl.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h"
#include "mlir/Dialect/NVGPU/Transforms/Transforms.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
#include "mlir/Dialect/SCF/Transforms/Transforms.h"
#include "mlir/Dialect/Utils/IndexingUtils.h"
#include "mlir/Dialect/Utils/StaticValueUtils.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/IR/AffineExpr.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Value.h"
#include "llvm/ADT/ArrayRef.h"
#include "mlir/Dialect/NVGPU/TransformOps/NVGPUTransformOps.cpp.inc"

Classes
struct	RowColIndexing
	Helper struct to encode a pair of row/column indexings in the form of affine expressions. More...
struct	MmaSyncBuilder
	Helper struct to provide a simple mapping from matmul operations to the corresponding mma.sync operation. More...
struct	HopperBuilder
	Helper to create the base Hopper-specific operations that are reused in various other places. More...
struct	CopyBuilder
	Helper to create the tma operations corresponding to linalg::CopyOp. More...

Macros
#define	DEBUG_TYPE "nvgpu-transforms"
#define	GET_OP_LIST
#define	MLIR_IR_BUILDERS_H
#define	MLIR_IR_OPDEFINITION_H
#define	MLIR_IR_DIALECT_H
#define	MLIR_IR_DIALECTREGISTRY_H
#define	MLIR_IR_MLIRCONTEXT_H
#define	MLIR_SUPPORT_LLVM_H
#define	MLIR_SUPPORT_TYPEID_H
#define	MLIR_DECLARE_EXPLICIT_SELF_OWNING_TYPE_ID(CLASS_NAME)
#define	MLIR_DEFINE_EXPLICIT_SELF_OWNING_TYPE_ID(CLASS_NAME)
#define	MLIR_DECLARE_EXPLICIT_FALLBACK_TYPE_ID(CLASS_NAME)
#define	MLIR_DEFINE_EXPLICIT_FALLBACK_TYPE_ID(CLASS_NAME)
#define	MLIR_USE_FALLBACK_TYPE_IDS false
#define	MLIR_DECLARE_EXPLICIT_TYPE_ID(CLASS_NAME)
#define	MLIR_DEFINE_EXPLICIT_TYPE_ID(CLASS_NAME)
#define	MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(CLASS_NAME)
#define	MLIR_IR_OPERATIONSUPPORT_H
#define	MLIR_IR_ATTRIBUTES_H
#define	MLIR_IR_ATTRIBUTESUPPORT_H
#define	MLIR_IR_STORAGEUNIQUERSUPPORT_H
#define	MLIR_IR_ATTRTYPESUBELEMENTS_H
#define	MLIR_IR_VISITORS_H
#define	MLIR_SUPPORT_WALKRESULT_H
#define	MLIR_SUPPORT_CYCLICREPLACERCACHE_H
#define	MLIR_SUPPORT_INTERFACESUPPORT_H
#define	MLIR_SUPPORT_STORAGEUNIQUER_H
#define	MLIR_IR_TYPES_H
#define	MLIR_IR_TYPESUPPORT_H
#define	MLIR_IR_BLOCKSUPPORT_H
#define	MLIR_IR_VALUE_H
#define	MLIR_IR_USEDEFLISTS_H
#define	MLIR_IR_LOCATION_H
#define	GET_ATTRDEF_CLASSES
#define	MLIR_IR_BUILTINATTRIBUTES_H
#define	MLIR_IR_BUILTINATTRIBUTEINTERFACES_H
#define	MLIR_IR_AFFINEMAP_H
#define	MLIR_IR_AFFINEEXPR_H
#define	MLIR_IR_BUILTINTYPEINTERFACES_H
#define	MLIR_IR_OPASMSUPPORT_H_
#define	MLIR_IR_BLOCK_H
#define	MLIR_SUPPORT_COMPLEX_H
#define	GET_ATTRDEF_CLASSES
#define	MLIR_IR_DIAGNOSTICS_H
#define	MLIR_IR_TYPERANGE_H
#define	MLIR_IR_VALUERANGE_H
#define	MLIR_IR_ODSSUPPORT_H
#define	MLIR_IR_OPERATION_H
#define	MLIR_IR_REGION_H
#define	MLIR_IR_BUILTINTYPES_H
#define	MLIR_SUPPORT_ADTEXTRAS_H
#define	MLIR_IR_QuantStorageTypeInterface_H
#define	GET_TYPEDEF_CLASSES
#define	MLIR_IR_OPIMPLEMENTATION_H
#define	MLIR_IR_DIALECTINTERFACE_H
#define	GET_OP_CLASSES

Functions
static bool	hasDefaultMemorySpace (BaseMemRefType type)
	Returns true if the given type has the default memory space.
static bool	hasSharedMemorySpace (BaseMemRefType type)
	Returns true if the given type has the shared (workgroup) memory space.
static Value	getValueLoadedFromGlobal (Operation *op)
	Returns the value produced by a load from the default memory space.
static bool	isStoreToShared (Operation *op, Value v)
	Returns true if the operation is storing the given value into shared memory.
static bool	isLoadFromGlobalStoredToShared (Operation *op)
	Returns true if the operation is a load from the default memory space the result of which is only stored into the shared memory space.
static LogicalResult	collectStage0PipeliningOps (scf::ForOp forOp, llvm::SmallPtrSet< Operation *, 16 > &ops)
	Populate ops with the set of operations that belong to the stage 0 of the pipelined version of the given loop when pipelining copies to shared memory.
static void	setAsyncWaitGroupsInFlight (OpBuilder &builder, Operation *op, scf::PipeliningOption::PipelinerPart part, unsigned iteration, unsigned depth)
	Hook for the loop pipeliner that sets the "num groups in flight" attribute of async wait operations corresponding to pipelined shared memory copies.
static void	getPipelineStages (scf::ForOp forOp, std::vector< std::pair< Operation , unsigned > > &opsWithPipelineStages, unsigned depth, llvm::SmallPtrSetImpl< Operation > &stage0Ops)
	Hook for the loop pipeliner that populates ops with the stage information as follows:
static Operation *	replaceOpWithPredicatedOp (RewriterBase &rewriter, Operation *op, Value predicate)
	Hook for the loop pipeliner.
static std::tuple< DiagnosedSilenceableFailure, scf::ForOp >	pipelineForSharedCopies (RewriterBase &rewriter, scf::ForOp forOp, int64_t depth, bool epiloguePeeling)
	Applies loop pipelining with the given depth to the given loop so that copies into the shared memory are pipelined.
template<typename ApplyFn, typename ReduceFn>
static void	foreachIndividualVectorElement (Value vector, ApplyFn applyFn, ReduceFn reduceFn)
	Helper functions to create customizable load and stores operations.
static std::tuple< SmallVector< int64_t >, SmallVector< int64_t >, SmallVector< int64_t > >	makeVectorShapes (ArrayRef< int64_t > lhs, ArrayRef< int64_t > rhs, ArrayRef< int64_t > res)
static Attribute	getSharedAddressSpaceAttribute (OpBuilder &b)

Macro Definition Documentation

◆ DEBUG_TYPE

#define DEBUG_TYPE "nvgpu-transforms"

Definition at line 40 of file NVGPUTransformOps.cpp.

◆ GET_ATTRDEF_CLASSES [1/2]

#define GET_ATTRDEF_CLASSES

◆ GET_ATTRDEF_CLASSES [2/2]

#define GET_ATTRDEF_CLASSES

◆ GET_OP_CLASSES

#define GET_OP_CLASSES

Definition at line 1137 of file NVGPUTransformOps.cpp.

◆ GET_OP_LIST

#define GET_OP_LIST

◆ GET_TYPEDEF_CLASSES

#define GET_TYPEDEF_CLASSES

◆ MLIR_DECLARE_EXPLICIT_FALLBACK_TYPE_ID

#define MLIR_DECLARE_EXPLICIT_FALLBACK_TYPE_ID ( CLASS_NAME )

Value:

  namespace mlir {                                                             \
  namespace detail {                                                           \
  template <>                                                                  \
  class TypeIDResolver<CLASS_NAME> : public FallbackTypeIDResolver {           \
  public:                                                                      \
    static TypeID resolveTypeID() {                                            \
      static_assert(is_fully_resolved<CLASS_NAME>(),                           \
                    "TypeID::get<> requires the complete definition of `T`");  \
      static TypeID id =                                                       \
          registerImplicitTypeID(llvm::getTypeName<CLASS_NAME>());             \
      return id;                                                               \
    }                                                                          \
  };                                                                           \
  } /* namespace detail */                                                     \
  } /* namespace mlir */

◆ MLIR_DECLARE_EXPLICIT_SELF_OWNING_TYPE_ID

#define MLIR_DECLARE_EXPLICIT_SELF_OWNING_TYPE_ID ( CLASS_NAME )

Value:

  namespace mlir {                                                             \
  namespace detail {                                                           \
  template <>                                                                  \
  class TypeIDResolver<CLASS_NAME> {                                           \
  public:                                                                      \
    static TypeID resolveTypeID() { return id; }                               \
                                                                               \
  private:                                                                     \
    static SelfOwningTypeID id;                                                \
  };                                                                           \
  } /* namespace detail */                                                     \
  } /* namespace mlir */

◆ MLIR_DECLARE_EXPLICIT_TYPE_ID

#define MLIR_DECLARE_EXPLICIT_TYPE_ID ( CLASS_NAME )

Value:

MLIR_DECLARE_EXPLICIT_SELF_OWNING_TYPE_ID(CLASS_NAME)

MLIR_DECLARE_EXPLICIT_SELF_OWNING_TYPE_ID

#define MLIR_DECLARE_EXPLICIT_SELF_OWNING_TYPE_ID(CLASS_NAME)

Definition TypeID.h:262

◆ MLIR_DEFINE_EXPLICIT_FALLBACK_TYPE_ID

#define MLIR_DEFINE_EXPLICIT_FALLBACK_TYPE_ID ( CLASS_NAME )

◆ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID

#define MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID ( CLASS_NAME )

Value:

  static ::mlir::TypeID resolveTypeID() {                                      \
    static ::mlir::SelfOwningTypeID id;                                        \
    return id;                                                                 \
  }                                                                            \
  static_assert(                                                               \
      ::mlir::detail::InlineTypeIDResolver::has_resolve_typeid<                \
          CLASS_NAME>::value,                                                  \
      "`MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID` must be placed in a "    \
      "public section of `" #CLASS_NAME "`");

◆ MLIR_DEFINE_EXPLICIT_SELF_OWNING_TYPE_ID

#define MLIR_DEFINE_EXPLICIT_SELF_OWNING_TYPE_ID ( CLASS_NAME )

Value:

  namespace mlir {                                                             \
  namespace detail {                                                           \
  SelfOwningTypeID TypeIDResolver<CLASS_NAME>::id = {};                        \
  } /* namespace detail */                                                     \
  } /* namespace mlir */

◆ MLIR_DEFINE_EXPLICIT_TYPE_ID

#define MLIR_DEFINE_EXPLICIT_TYPE_ID ( CLASS_NAME )

Value:

MLIR_DEFINE_EXPLICIT_SELF_OWNING_TYPE_ID(CLASS_NAME)

MLIR_DEFINE_EXPLICIT_SELF_OWNING_TYPE_ID

#define MLIR_DEFINE_EXPLICIT_SELF_OWNING_TYPE_ID(CLASS_NAME)

Definition TypeID.h:276

◆ MLIR_IR_AFFINEEXPR_H

#define MLIR_IR_AFFINEEXPR_H

◆ MLIR_IR_AFFINEMAP_H

#define MLIR_IR_AFFINEMAP_H

◆ MLIR_IR_ATTRIBUTES_H

#define MLIR_IR_ATTRIBUTES_H

◆ MLIR_IR_ATTRIBUTESUPPORT_H

#define MLIR_IR_ATTRIBUTESUPPORT_H

◆ MLIR_IR_ATTRTYPESUBELEMENTS_H

#define MLIR_IR_ATTRTYPESUBELEMENTS_H

◆ MLIR_IR_BLOCK_H

#define MLIR_IR_BLOCK_H

◆ MLIR_IR_BLOCKSUPPORT_H

#define MLIR_IR_BLOCKSUPPORT_H

◆ MLIR_IR_BUILDERS_H

#define MLIR_IR_BUILDERS_H

◆ MLIR_IR_BUILTINATTRIBUTEINTERFACES_H

#define MLIR_IR_BUILTINATTRIBUTEINTERFACES_H

◆ MLIR_IR_BUILTINATTRIBUTES_H

#define MLIR_IR_BUILTINATTRIBUTES_H

◆ MLIR_IR_BUILTINTYPEINTERFACES_H

#define MLIR_IR_BUILTINTYPEINTERFACES_H

◆ MLIR_IR_BUILTINTYPES_H

#define MLIR_IR_BUILTINTYPES_H

◆ MLIR_IR_DIAGNOSTICS_H

#define MLIR_IR_DIAGNOSTICS_H

◆ MLIR_IR_DIALECT_H

#define MLIR_IR_DIALECT_H

◆ MLIR_IR_DIALECTINTERFACE_H

#define MLIR_IR_DIALECTINTERFACE_H

◆ MLIR_IR_DIALECTREGISTRY_H

#define MLIR_IR_DIALECTREGISTRY_H

◆ MLIR_IR_LOCATION_H

#define MLIR_IR_LOCATION_H

◆ MLIR_IR_MLIRCONTEXT_H

#define MLIR_IR_MLIRCONTEXT_H

◆ MLIR_IR_ODSSUPPORT_H

#define MLIR_IR_ODSSUPPORT_H

◆ MLIR_IR_OPASMSUPPORT_H_

#define MLIR_IR_OPASMSUPPORT_H_

◆ MLIR_IR_OPDEFINITION_H

#define MLIR_IR_OPDEFINITION_H

◆ MLIR_IR_OPERATION_H

#define MLIR_IR_OPERATION_H

◆ MLIR_IR_OPERATIONSUPPORT_H

#define MLIR_IR_OPERATIONSUPPORT_H

◆ MLIR_IR_OPIMPLEMENTATION_H

#define MLIR_IR_OPIMPLEMENTATION_H

◆ MLIR_IR_QuantStorageTypeInterface_H

#define MLIR_IR_QuantStorageTypeInterface_H

◆ MLIR_IR_REGION_H

#define MLIR_IR_REGION_H

◆ MLIR_IR_STORAGEUNIQUERSUPPORT_H

#define MLIR_IR_STORAGEUNIQUERSUPPORT_H

◆ MLIR_IR_TYPERANGE_H

#define MLIR_IR_TYPERANGE_H

◆ MLIR_IR_TYPES_H

#define MLIR_IR_TYPES_H

◆ MLIR_IR_TYPESUPPORT_H

#define MLIR_IR_TYPESUPPORT_H

◆ MLIR_IR_USEDEFLISTS_H

#define MLIR_IR_USEDEFLISTS_H

◆ MLIR_IR_VALUE_H

#define MLIR_IR_VALUE_H

◆ MLIR_IR_VALUERANGE_H

#define MLIR_IR_VALUERANGE_H

◆ MLIR_IR_VISITORS_H

#define MLIR_IR_VISITORS_H

◆ MLIR_SUPPORT_ADTEXTRAS_H

#define MLIR_SUPPORT_ADTEXTRAS_H

◆ MLIR_SUPPORT_COMPLEX_H

#define MLIR_SUPPORT_COMPLEX_H

◆ MLIR_SUPPORT_CYCLICREPLACERCACHE_H

#define MLIR_SUPPORT_CYCLICREPLACERCACHE_H

◆ MLIR_SUPPORT_INTERFACESUPPORT_H

#define MLIR_SUPPORT_INTERFACESUPPORT_H

◆ MLIR_SUPPORT_LLVM_H

#define MLIR_SUPPORT_LLVM_H

◆ MLIR_SUPPORT_STORAGEUNIQUER_H

#define MLIR_SUPPORT_STORAGEUNIQUER_H

◆ MLIR_SUPPORT_TYPEID_H

#define MLIR_SUPPORT_TYPEID_H

◆ MLIR_SUPPORT_WALKRESULT_H

#define MLIR_SUPPORT_WALKRESULT_H

◆ MLIR_USE_FALLBACK_TYPE_IDS

#define MLIR_USE_FALLBACK_TYPE_IDS false

Function Documentation

◆ collectStage0PipeliningOps()

LogicalResult collectStage0PipeliningOps	(	scf::ForOp	forOp,
		llvm::SmallPtrSet< Operation *, 16 > &	ops )

static

Populate ops with the set of operations that belong to the stage 0 of the pipelined version of the given loop when pipelining copies to shared memory.

Specifically, this collects:

all loads from global memory, both sync and async;
the barriers for async loads.

In particular, barriers are omitted if they do not dominate at least one async load for which there is not yet a barrier.

Definition at line 192 of file NVGPUTransformOps.cpp.

References isLoadFromGlobalStoredToShared(), and success().

Referenced by pipelineForSharedCopies().

◆ foreachIndividualVectorElement()

template<typename ApplyFn, typename ReduceFn>

void foreachIndividualVectorElement	(	Value	vector,
		ApplyFn	applyFn,
		ReduceFn	reduceFn )

static

Helper functions to create customizable load and stores operations.

The specific shapes of each MMA instruction are passed via the IndexCalculator callback.

Definition at line 634 of file NVGPUTransformOps.cpp.

References mlir::computeStrides(), mlir::delinearize(), indices, and vectorShape().

◆ getPipelineStages()

void getPipelineStages	(	scf::ForOp	forOp,
		std::vector< std::pair< Operation *, unsigned > > &	opsWithPipelineStages,
		unsigned	depth,
		llvm::SmallPtrSetImpl< Operation * > &	stage0Ops )

static

Hook for the loop pipeliner that populates ops with the stage information as follows:

operations in stage0Ops (typically loads from global memory and related barriers) are at stage 0;
operations in the backward slice of any stage0Ops are all at stage 0;
other operations are at stage depth;
the internal order of the pipelined loop has ops at stage depth first, then those at stage 0, with relative order within each group preserved.

Definition at line 263 of file NVGPUTransformOps.cpp.

References mlir::getBackwardSlice(), mlir::Operation::getBlock(), options, and result.

Referenced by pipelineForSharedCopies().

◆ getSharedAddressSpaceAttribute()

Attribute getSharedAddressSpaceAttribute ( OpBuilder & b )

static

Definition at line 910 of file NVGPUTransformOps.cpp.

References b.

Referenced by HopperBuilder::buildAndInitBarrierInSharedMemory(), and HopperBuilder::buildGlobalMemRefDescriptor().

◆ getValueLoadedFromGlobal()

Value getValueLoadedFromGlobal ( Operation * op )

static

Returns the value produced by a load from the default memory space.

Returns null if the operation is not such a load.

Definition at line 149 of file NVGPUTransformOps.cpp.

References hasDefaultMemorySpace(), and load.

Referenced by isLoadFromGlobalStoredToShared().

◆ hasDefaultMemorySpace()

bool hasDefaultMemorySpace ( BaseMemRefType type )

static

Returns true if the given type has the default memory space.

Definition at line 135 of file NVGPUTransformOps.cpp.

References mlir::BaseMemRefType::getMemorySpace(), and mlir::BaseMemRefType::getMemorySpaceAsInt().

Referenced by getValueLoadedFromGlobal().

◆ hasSharedMemorySpace()

bool hasSharedMemorySpace ( BaseMemRefType type )

static

Returns true if the given type has the shared (workgroup) memory space.

Definition at line 140 of file NVGPUTransformOps.cpp.

References mlir::BaseMemRefType::getMemorySpace().

Referenced by isStoreToShared().

◆ isLoadFromGlobalStoredToShared()

bool isLoadFromGlobalStoredToShared ( Operation * op )

static

Returns true if the operation is a load from the default memory space the result of which is only stored into the shared memory space.

Definition at line 174 of file NVGPUTransformOps.cpp.

References mlir::Value::getUsers(), getValueLoadedFromGlobal(), mlir::Value::hasOneUse(), and isStoreToShared().

Referenced by collectStage0PipeliningOps().

◆ isStoreToShared()

bool isStoreToShared	(	Operation *	op,
		Value	v )

static

Returns true if the operation is storing the given value into shared memory.

Definition at line 162 of file NVGPUTransformOps.cpp.

References hasSharedMemorySpace().

Referenced by isLoadFromGlobalStoredToShared().

◆ makeVectorShapes()

std::tuple< SmallVector< int64_t >, SmallVector< int64_t >, SmallVector< int64_t > > makeVectorShapes	(	ArrayRef< int64_t >	lhs,
		ArrayRef< int64_t >	rhs,
		ArrayRef< int64_t >	res )

static

Definition at line 723 of file NVGPUTransformOps.cpp.

◆ pipelineForSharedCopies()

std::tuple< DiagnosedSilenceableFailure, scf::ForOp > pipelineForSharedCopies	(	RewriterBase &	rewriter,
		scf::ForOp	forOp,
		int64_t	depth,
		bool	epiloguePeeling )

static

Applies loop pipelining with the given depth to the given loop so that copies into the shared memory are pipelined.

Doesn't affect other loops. Returns a pair containing the error state and the pipelined op, the latter being null in case of any failure. The error state contains a definite error if the IR has been modified and a silenceable error otherwise.

Definition at line 337 of file NVGPUTransformOps.cpp.

References collectStage0PipeliningOps(), mlir::DiagnosedSilenceableFailure::definiteFailure(), mlir::emitSilenceableFailure(), getPipelineStages(), options, replaceOpWithPredicatedOp(), setAsyncWaitGroupsInFlight(), mlir::OpBuilder::setInsertionPoint(), and mlir::DiagnosedSilenceableFailure::success().

◆ replaceOpWithPredicatedOp()

Operation * replaceOpWithPredicatedOp	(	RewriterBase &	rewriter,
		Operation *	op,
		Value	predicate )

static

Hook for the loop pipeliner.

Replaces op with a predicated version and returns the resulting operation. Returns the original op if the predication isn't necessary for the given op. Returns null if predication is needed but not supported.

Definition at line 294 of file NVGPUTransformOps.cpp.

References mlir::arith::ConstantIndexOp::create(), mlir::isMemoryEffectFree(), and mlir::RewriterBase::replaceOp().

Referenced by pipelineForSharedCopies().

◆ setAsyncWaitGroupsInFlight()

void setAsyncWaitGroupsInFlight	(	OpBuilder &	builder,
		Operation *	op,
		scf::PipeliningOption::PipelinerPart	part,
		unsigned	iteration,
		unsigned	depth )

static

Hook for the loop pipeliner that sets the "num groups in flight" attribute of async wait operations corresponding to pipelined shared memory copies.

Definition at line 229 of file NVGPUTransformOps.cpp.

References mlir::scf::PipeliningOption::Epilogue, mlir::scf::PipeliningOption::Kernel, and mlir::scf::PipeliningOption::Prologue.

Referenced by pipelineForSharedCopies().

Classes

Macros

Functions

Macro Definition Documentation

◆ DEBUG_TYPE

◆ GET_ATTRDEF_CLASSES [1/2]

◆ GET_ATTRDEF_CLASSES [2/2]

◆ GET_OP_CLASSES

◆ GET_OP_LIST

◆ GET_TYPEDEF_CLASSES

◆ MLIR_DECLARE_EXPLICIT_FALLBACK_TYPE_ID

◆ MLIR_DECLARE_EXPLICIT_SELF_OWNING_TYPE_ID

◆ MLIR_DECLARE_EXPLICIT_TYPE_ID

◆ MLIR_DEFINE_EXPLICIT_FALLBACK_TYPE_ID

◆ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID

◆ MLIR_DEFINE_EXPLICIT_SELF_OWNING_TYPE_ID

◆ MLIR_DEFINE_EXPLICIT_TYPE_ID

◆ MLIR_IR_AFFINEEXPR_H

◆ MLIR_IR_AFFINEMAP_H

◆ MLIR_IR_ATTRIBUTES_H

◆ MLIR_IR_ATTRIBUTESUPPORT_H

◆ MLIR_IR_ATTRTYPESUBELEMENTS_H

◆ MLIR_IR_BLOCK_H

◆ MLIR_IR_BLOCKSUPPORT_H

◆ MLIR_IR_BUILDERS_H

◆ MLIR_IR_BUILTINATTRIBUTEINTERFACES_H

◆ MLIR_IR_BUILTINATTRIBUTES_H

◆ MLIR_IR_BUILTINTYPEINTERFACES_H

◆ MLIR_IR_BUILTINTYPES_H

◆ MLIR_IR_DIAGNOSTICS_H

◆ MLIR_IR_DIALECT_H

◆ MLIR_IR_DIALECTINTERFACE_H

◆ MLIR_IR_DIALECTREGISTRY_H

◆ MLIR_IR_LOCATION_H

◆ MLIR_IR_MLIRCONTEXT_H

◆ MLIR_IR_ODSSUPPORT_H

◆ MLIR_IR_OPASMSUPPORT_H_

◆ MLIR_IR_OPDEFINITION_H

◆ MLIR_IR_OPERATION_H

◆ MLIR_IR_OPERATIONSUPPORT_H

◆ MLIR_IR_OPIMPLEMENTATION_H

◆ MLIR_IR_QuantStorageTypeInterface_H

◆ MLIR_IR_REGION_H

◆ MLIR_IR_STORAGEUNIQUERSUPPORT_H

◆ MLIR_IR_TYPERANGE_H

◆ MLIR_IR_TYPES_H

◆ MLIR_IR_TYPESUPPORT_H

◆ MLIR_IR_USEDEFLISTS_H

◆ MLIR_IR_VALUE_H

◆ MLIR_IR_VALUERANGE_H

◆ MLIR_IR_VISITORS_H

◆ MLIR_SUPPORT_ADTEXTRAS_H

◆ MLIR_SUPPORT_COMPLEX_H

◆ MLIR_SUPPORT_CYCLICREPLACERCACHE_H

◆ MLIR_SUPPORT_INTERFACESUPPORT_H

◆ MLIR_SUPPORT_LLVM_H

◆ MLIR_SUPPORT_STORAGEUNIQUER_H

◆ MLIR_SUPPORT_TYPEID_H

◆ MLIR_SUPPORT_WALKRESULT_H

◆ MLIR_USE_FALLBACK_TYPE_IDS

Function Documentation

◆ collectStage0PipeliningOps()

◆ foreachIndividualVectorElement()

◆ getPipelineStages()

◆ getSharedAddressSpaceAttribute()

◆ getValueLoadedFromGlobal()

◆ hasDefaultMemorySpace()

◆ hasSharedMemorySpace()

◆ isLoadFromGlobalStoredToShared()

◆ isStoreToShared()

◆ makeVectorShapes()

◆ pipelineForSharedCopies()

◆ replaceOpWithPredicatedOp()

◆ setAsyncWaitGroupsInFlight()