Home
last modified time | relevance | path

Searched refs:srcMemref (Results 1 – 7 of 7) sorted by relevance

/llvm-project-15.0.7/mlir/lib/Dialect/NVGPU/IR/
H A DNVGPUDialect.cpp69 auto srcMemref = getSrc().getType().cast<MemRefType>(); in verify() local
72 if (!isLastMemrefDimUnitStride(srcMemref)) in verify()
79 if (dstMemref.getElementType() != srcMemref.getElementType()) in verify()
81 if (size_t(srcMemref.getRank()) != getSrcIndices().size()) in verify()
82 return emitOpError() << "expected " << srcMemref.getRank() in verify()
199 auto srcMemref = getSrcMemref().getType().cast<MemRefType>(); in verify() local
225 if (!(srcMemref.getMemorySpaceAsInt() == smemAddressSpace)) in verify()
/llvm-project-15.0.7/mlir/lib/Dialect/Tensor/Transforms/
H A DBufferizableOpInterfaceImpl.cpp287 FailureOr<Value> srcMemref = in bufferize() local
289 if (failed(srcMemref)) in bufferize()
291 auto srcMemrefType = srcMemref->getType().cast<MemRefType>(); in bufferize()
300 loc, subviewMemRefType, *srcMemref, mixedOffsets, mixedSizes, in bufferize()
330 FailureOr<Value> srcMemref = in bufferize() local
332 if (failed(srcMemref)) in bufferize()
334 replaceOpWithNewBufferizedOp<memref::LoadOp>(rewriter, op, *srcMemref, in bufferize()
730 FailureOr<Value> srcMemref = in bufferize() local
732 if (failed(srcMemref)) in bufferize()
734 if (failed(options.createMemCpy(rewriter, loc, *srcMemref, subView))) in bufferize()
/llvm-project-15.0.7/mlir/include/mlir/Dialect/NVGPU/IR/
H A DNVGPU.td58 PredOpTrait<"srcMemref and res have same element type",
77 let arguments = (ins Arg<AnyMemRef, "", [MemRead]>:$srcMemref,
82 $srcMemref`[` $indices `]` attr-dict `:` type($srcMemref) `->` type($res)
/llvm-project-15.0.7/mlir/test/Dialect/NVGPU/
H A Dinvalid.mlir5 // expected-error @+1 {{expected nvgpu.ldmatrix srcMemref must have memory space 3}}
45 …// expected-error @+1 {{'nvgpu.ldmatrix' op failed to verify that srcMemref and res have same elem…
/llvm-project-15.0.7/mlir/lib/Conversion/GPUToNVVM/
H A DWmmaOpsToNvvm.cpp115 loc, subgroupMmaLoadMatrixOp.srcMemref().getType().cast<MemRefType>(), in matchAndRewrite()
116 adaptor.srcMemref(), adaptor.indices(), rewriter); in matchAndRewrite()
/llvm-project-15.0.7/mlir/include/mlir/Dialect/GPU/IR/
H A DGPUOps.td1095 let arguments = (ins Arg<MemRefOf<[F16, F32]>, "", [MemRead]>:$srcMemref,
1102 $srcMemref`[`$indices`]` attr-dict `:` type($srcMemref) `->` type($res)
/llvm-project-15.0.7/mlir/lib/Dialect/GPU/IR/
H A DGPUDialect.cpp1167 auto srcType = srcMemref().getType(); in verify()