Home
last modified time | relevance | path

Searched refs:MMAMatrixType (Results 1 – 9 of 9) sorted by relevance

/llvm-project-15.0.7/mlir/include/mlir/Dialect/GPU/IR/
H A DGPUDialect.h123 class MMAMatrixType
124 : public Type::TypeBase<MMAMatrixType, Type, MMAMatrixStorageType> {
129 static MMAMatrixType get(ArrayRef<int64_t> shape, Type elementType,
134 static MMAMatrixType getChecked(function_ref<InFlightDiagnostic()> emitError,
H A DGPUBase.td69 // Predicat to check if type is gpu::MMAMatrixType.
70 def IsMMAMatrixTypePred : CPred<"$_self.isa<::mlir::gpu::MMAMatrixType>()">;
77 "$_self.cast<::mlir::gpu::MMAMatrixType>().getElementType()",
78 "gpu.mma_matrix", "::mlir::gpu::MMAMatrixType">;
H A DGPUOps.td1186 "$_self.cast<gpu::MMAMatrixType>().getElementType()">]>{
1217 gpu::MMAMatrixType getType() {
1218 return res().getType().cast<gpu::MMAMatrixType>();
1278 gpu::MMAMatrixType getType() {
1279 return res().getType().cast<gpu::MMAMatrixType>();
/llvm-project-15.0.7/mlir/lib/Conversion/GPUToNVVM/
H A DWmmaOpsToNvvm.cpp54 static NVVM::MMATypes getElementType(gpu::MMAMatrixType type) { in getElementType()
82 gpu::MMAMatrixType retType = in matchAndRewrite()
83 subgroupMmaLoadMatrixOp.res().getType().cast<gpu::MMAMatrixType>(); in matchAndRewrite()
149 gpu::MMAMatrixType srcType = in matchAndRewrite()
214 gpu::MMAMatrixType aType = in matchAndRewrite()
215 subgroupMmaComputeOp.opA().getType().cast<gpu::MMAMatrixType>(); in matchAndRewrite()
217 gpu::MMAMatrixType cType = in matchAndRewrite()
218 subgroupMmaComputeOp.opC().getType().cast<gpu::MMAMatrixType>(); in matchAndRewrite()
257 subgroupMmaConstantOp.getType().cast<gpu::MMAMatrixType>()); in matchAndRewrite()
335 subgroupMmaElementwiseOp.getType().cast<gpu::MMAMatrixType>()); in matchAndRewrite()
[all …]
H A DLowerGpuOpsToNVVMOps.cpp194 converter.addConversion([&](gpu::MMAMatrixType type) -> Type { in runOnOperation()
/llvm-project-15.0.7/mlir/lib/Dialect/GPU/IR/
H A DGPUDialect.cpp40 MMAMatrixType MMAMatrixType::get(ArrayRef<int64_t> shape, Type elementType, in get()
45 MMAMatrixType
55 ArrayRef<int64_t> MMAMatrixType::getShape() const { in getShape()
63 bool MMAMatrixType::isValidElementType(Type elementType) { in isValidElementType()
78 if (!MMAMatrixType::isValidElementType(elementType)) in verify()
121 addTypes<MMAMatrixType>(); in initialize()
183 .Case<MMAMatrixType>([&](MMAMatrixType fragTy) { in printType()
1169 auto resMatrixType = resType.cast<gpu::MMAMatrixType>(); in verify()
1198 auto srcMatrixType = srcType.cast<gpu::MMAMatrixType>(); in verify()
1224 SmallVector<MMAMatrixType, 3> opTypes; in verify()
[all …]
/llvm-project-15.0.7/mlir/include/mlir/Conversion/GPUToNVVM/
H A DGPUToNVVMPass.h25 class MMAMatrixType; variable
28 LLVM::LLVMStructType convertMMAToLLVMType(gpu::MMAMatrixType type);
/llvm-project-15.0.7/mlir/lib/Conversion/VectorToGPU/
H A DVectorToGPU.cpp407 gpu::MMAMatrixType type = in convertTransferReadOp()
408 gpu::MMAMatrixType::get(op.getVectorType().getShape(), in convertTransferReadOp()
707 gpu::MMAMatrixType type = gpu::MMAMatrixType::get( in convertConstantOp()
721 gpu::MMAMatrixType type = gpu::MMAMatrixType::get( in convertBroadcastOp()
/llvm-project-15.0.7/mlir/test/Dialect/GPU/
H A Dinvalid.mlir462 // expected-error @+1 {{MMAMatrixType must have exactly two dimensions}}
482 // expected-error @+1 {{MMAMatrixType elements must be F16 or F32}}