Home
last modified time | relevance | path

Searched refs:MMATypes (Results 1 – 5 of 5) sorted by relevance

/llvm-project-15.0.7/mlir/lib/Dialect/LLVMIR/IR/
H A DNVVMDialect.cpp82 return NVVM::MMATypes::f64; in inferOperandMMAType()
105 return (type == MMATypes::u4 || type == MMATypes::s4); in isInt4PtxType()
109 return (type == MMATypes::u8 || type == MMATypes::s8); in isInt8PtxType()
376 case MMATypes::tf32: in verify()
382 case MMATypes::f16: in verify()
383 case MMATypes::bf16: in verify()
389 case MMATypes::s4: in verify()
390 case MMATypes::u4: in verify()
393 case MMATypes::b1: in verify()
396 case MMATypes::s8: in verify()
[all …]
/llvm-project-15.0.7/mlir/lib/Conversion/NVGPUToNVVM/
H A DNVGPUToNVVM.cpp274 NVVM::MMATypes ptxTypeA; in matchAndRewrite()
275 NVVM::MMATypes ptxTypeB; in matchAndRewrite()
285 ptxTypeA = NVVM::MMATypes::s8; in matchAndRewrite()
286 ptxTypeB = NVVM::MMATypes::s8; in matchAndRewrite()
289 ptxTypeA = NVVM::MMATypes::s4; in matchAndRewrite()
290 ptxTypeB = NVVM::MMATypes::s4; in matchAndRewrite()
293 ptxTypeA = NVVM::MMATypes::f16; in matchAndRewrite()
294 ptxTypeB = NVVM::MMATypes::f16; in matchAndRewrite()
296 ptxTypeA = NVVM::MMATypes::f64; in matchAndRewrite()
297 ptxTypeB = NVVM::MMATypes::f64; in matchAndRewrite()
[all …]
/llvm-project-15.0.7/mlir/lib/Conversion/GPUToNVVM/
H A DWmmaOpsToNvvm.cpp54 static NVVM::MMATypes getElementType(gpu::MMAMatrixType type) { in getElementType()
56 return NVVM::MMATypes::f16; in getElementType()
58 return type.getOperand().equals("COp") ? NVVM::MMATypes::f32 in getElementType()
59 : NVVM::MMATypes::tf32; in getElementType()
88 NVVM::MMATypes eltype = getElementType(retType); in matchAndRewrite()
153 NVVM::MMATypes eltype = getElementType(srcType); in matchAndRewrite()
224 NVVM::MMATypes sourceType = getElementType(aType); in matchAndRewrite()
225 NVVM::MMATypes destType = getElementType(cType); in matchAndRewrite()
364 NVVM::MMATypes eltType = getElementType(type); in convertMMAToLLVMType()
/llvm-project-15.0.7/mlir/include/mlir/Dialect/LLVMIR/
H A DNVVMOps.td616 def MMATypes : I32EnumAttr<"MMATypes", "NVVM MMA types",
624 def MMATypesAttr : EnumAttr<NVVM_Dialect, MMATypes, "mma_type"> {
660 "mlir::NVVM::MMATypes eltypeEnum,mlir::NVVM::MMAFrag fragEnum) {"
730 "mlir::NVVM::MMATypes eltypeEnum) {"
787 "mlir::NVVM::MMALayout layoutBEnum, mlir::NVVM::MMATypes eltypeAEnum,"
788 "mlir::NVVM::MMATypes eltypeBEnum) {"
963 mlir::NVVM::MMATypes eltypeAEnum, mlir::NVVM::MMATypes eltypeBEnum,
964 mlir::NVVM::MMATypes eltypeCEnum, mlir::NVVM::MMATypes eltypeDEnum) {
979 MMATypes accumPtxType();
980 MMATypes resultPtxType();
[all …]
H A DNVVMDialect.h39 std::pair<mlir::Type, unsigned> inferMMAType(mlir::NVVM::MMATypes type,