| /llvm-project/mlir/lib/Dialect/NVGPU/IR/ |
| H A D | NVGPUDialect.cpp | 126 ::mlir::OperationState &odsState, Value matrixA, in build() argument
128 build(odsBuilder, odsState, matrixC.getType(), matrixA, matrixB, matrixC, in build()
133 ::mlir::OperationState &odsState, Value matrixA, in build() argument
136 build(odsBuilder, odsState, matrixC.getType(), matrixA, matrixB, matrixC, in build()
143 TypedValue<VectorType> matrixA, in verifyMmaSyncOp() argument
166 auto aVector = matrixA.getType(); in verifyMmaSyncOp()
207 return op->emitError() << "matrixA must be 2 dimensional vector"; in verifyMmaSyncOp()
279 ::mlir::OperationState &odsState, Value matrixA, in verify()
282 build(odsBuilder, odsState, matrixC.getType(), matrixA, matrixB, matrixC, in verify()
566 MemRefType matrixA in verify()
267 build(::mlir::OpBuilder & odsBuilder,::mlir::OperationState & odsState,Value matrixA,Value matrixB,Value matrixC,Value sparseMetadata,ArrayRef<int64_t> mmaShape) build() argument
554 MemRefType matrixA = getDescriptorA().getType().getTensor(); verify() local
[all...] |
| /llvm-project/llvm/test/CodeGen/Hexagon/ |
| H A D | float-amode.ll | 22 %matrixA = getelementptr inbounds %struct.matrix_params, ptr %params, i32 0, i32 2
23 %0 = load ptr, ptr %matrixA, align 4
45 %matrixA = getelementptr inbounds %struct.matrix_params, ptr %params, i32 0, i32 2
46 %0 = load ptr, ptr %matrixA, align 4
64 %matrixA = getelementptr inbounds %struct.matrix_params, ptr %params, i32 0, i32 2
65 %0 = load ptr, ptr %matrixA, align 4
|
| /llvm-project/mlir/include/mlir/Dialect/NVGPU/IR/ |
| H A D | NVGPU.td | 268 PredOpTrait<"matrixA and matrixB have same element type",
300 %res = nvgpu.mma.sync (%matrixA, %matrixB, %matrixC) {mmaShape = [16, 8, 16]} :
304 let arguments = (ins AnyVectorOfNonZeroRank:$matrixA,
313 OpBuilder<(ins "Value":$matrixA,
317 OpBuilder<(ins "Value":$matrixA,
325 `(` $matrixA`,` $matrixB`,` $matrixC `)` attr-dict
326 `:` `(` type($matrixA) `,` type($matrixB) `,` type($matrixC) `)` `->` type($res)
339 where operand A is "structured sparse". In this case, the `matrixA` operand
360 let arguments = (ins AnyVectorOfNonZeroRank:$matrixA,
372 OpBuilder<(ins "Value":$matrixA,
[all...] |
| /llvm-project/mlir/test/Integration/GPU/CUDA/sm90/ |
| H A D | gemm_pred_f32_f16_f16_128x128x128.mlir | 106 %matrixA:2 = gpu.alloc async [%token] () : memref<128x128xf16>
109 %1 = gpu.memcpy async [%token] %matrixA, %matrixAHost : memref<128x128xf16>, memref<128x128xf16>
111 %castA = memref.cast %matrixA : memref<128x128xf16> to memref<*xf16>
|
| H A D | gemm_f32_f16_f16_128x128x128.mlir | 106 %matrixA:2 = gpu.alloc async [%token] () : memref<128x128xf16>
109 %1 = gpu.memcpy async [%token] %matrixA, %matrixAHost : memref<128x128xf16>, memref<128x128xf16>
111 %castA = memref.cast %matrixA : memref<128x128xf16> to memref<*xf16>
|
| /llvm-project/mlir/test/Conversion/VectorToGPU/ |
| H A D | fold-arith-vector-to-mma-ops-mma-sync.mlir | 4 // FP16 input, F32 accumulation row-row-row (ldmatrix x4 for matrixA and ldmatrix x4 for matrixB)
|
| H A D | vector-to-mma-ops-mma-sync.mlir | 169 // FP16 row-row-row (ldmatrix x4 for matrixA and ldmatrix x2 for matrixB)
203 // FP16 row-row-row (ldmatrix x4 for matrixA and ldmatrix x4 for matrixB)
|
| /llvm-project/flang/unittests/Optimizer/Builder/Runtime/ |
| H A D | TransformationalTest.cpp | 159 mlir::Value matrixA = builder.create<fir::UndefOp>(loc, boxTy1); in testGenMatmul() local
161 fir::runtime::genMatmul(builder, loc, result, matrixA, matrixB); in testGenMatmul()
|
| /llvm-project/mlir/test/Dialect/NVGPU/ |
| H A D | invalid.mlir | 130 // expected-error @+1 {{op failed to verify that matrixA and matrixB have same element type}}
361 // expected-error @+1 {{matrixA must be 2 dimensional vector}}
|
| /llvm-project/flang/lib/Optimizer/Builder/ |
| H A D | IntrinsicCall.cpp | 6035 mlir::Value matrixA = fir::getBase(matrixTmpA); in genScan()
6048 fir::runtime::genMatmul(builder, loc, resultIrBox, matrixA, matrixB); in genScan()
6062 mlir::Value matrixA = fir::getBase(matrixTmpA); in genScan()
6075 fir::runtime::genMatmulTranspose(builder, loc, resultIrBox, matrixA, matrixB); in genScan()
5324 mlir::Value matrixA = fir::getBase(matrixTmpA); genMatmul() local
5351 mlir::Value matrixA = fir::getBase(matrixTmpA); genMatmulTranspose() local
|