| /llvm-project-15.0.7/mlir/lib/Dialect/GPU/Transforms/ |
| H A D | MemoryPromotion.cpp | 42 if (rank < GPUDialect::getNumWorkgroupDimensions()) { in insertCopyLoops()
43 unsigned extraLoops = GPUDialect::getNumWorkgroupDimensions() - rank; in insertCopyLoops()
81 GPUDialect::getNumWorkgroupDimensions())))) { in insertCopyLoops()
150 int workgroupMemoryAddressSpace = gpu::GPUDialect::getWorkgroupAddressSpace(); in promoteToWorkgroupMemory()
|
| H A D | KernelOutlining.cpp | 168 outlinedFunc->setAttr(gpu::GPUDialect::getKernelFuncAttrName(), in outlineKernelFuncImpl()
329 getOperation()->setAttr(gpu::GPUDialect::getContainerModuleAttrName(), in runOnOperation()
|
| H A D | AllReduceLowering.cpp | 162 gpu::GPUDialect::getWorkgroupAddressSpace(); in createWorkgroupBuffer()
|
| /llvm-project-15.0.7/mlir/include/mlir/Dialect/GPU/Transforms/ |
| H A D | Passes.td | 17 let dependentDialects = ["mlir::gpu::GPUDialect"];
37 let dependentDialects = ["mlir::gpu::GPUDialect"];
|
| /llvm-project-15.0.7/mlir/lib/Dialect/GPU/IR/ |
| H A D | GPUDialect.cpp | 100 bool GPUDialect::isKernel(Operation *op) { in isKernel()
119 void GPUDialect::initialize() { in initialize()
133 Type GPUDialect::parseType(DialectAsmParser &parser) const { in parseType()
180 void GPUDialect::printType(Type type, DialectAsmPrinter &os) const { in printType()
194 LogicalResult GPUDialect::verifyOperationAttribute(Operation *op, in verifyOperationAttribute()
242 GPUDialect::getKernelFuncAttrName())) in verifyOperationAttribute()
244 << GPUDialect::getKernelFuncAttrName() << "' attribute"; in verifyOperationAttribute()
741 GPUDialect::getContainerModuleAttrName())) in verify()
743 GPUDialect::getContainerModuleAttrName() + in verify()
924 result.addAttribute(GPUDialect::getKernelFuncAttrName(), in parse()
[all …]
|
| /llvm-project-15.0.7/mlir/lib/CAPI/Dialect/ |
| H A D | GPU.cpp | 13 MLIR_DEFINE_CAPI_DIALECT_REGISTRATION(GPU, gpu, mlir::gpu::GPUDialect)
|
| /llvm-project-15.0.7/mlir/lib/Dialect/NVGPU/IR/ |
| H A D | NVGPUDialect.cpp | 71 unsigned workgroupAddressSpace = gpu::GPUDialect::getWorkgroupAddressSpace(); in verify()
219 unsigned smemAddressSpace = gpu::GPUDialect::getWorkgroupAddressSpace(); in verify()
|
| /llvm-project-15.0.7/mlir/test/lib/Dialect/SPIRV/ |
| H A D | TestEntryPointAbi.cpp | 56 if (!gpu::GPUDialect::isKernel(gpuFunc) || gpuFunc->getAttr(attrName)) in runOnOperation()
|
| /llvm-project-15.0.7/mlir/lib/Conversion/ |
| H A D | PassDetail.h | 42 class GPUDialect; variable
|
| /llvm-project-15.0.7/utils/bazel/llvm-project-overlay/mlir/test/ |
| H A D | BUILD.bazel | 442 "//mlir:GPUDialect",
470 "//mlir:GPUDialect",
564 "//mlir:GPUDialect",
614 "//mlir:GPUDialect",
|
| /llvm-project-15.0.7/mlir/tools/mlir-vulkan-runner/ |
| H A D | mlir-vulkan-runner.cpp | 77 mlir::gpu::GPUDialect, mlir::spirv::SPIRVDialect, in main()
|
| /llvm-project-15.0.7/mlir/include/mlir/ |
| H A D | InitAllDialects.h | 87 gpu::GPUDialect, in registerAllDialects()
|
| /llvm-project-15.0.7/mlir/tools/mlir-spirv-cpu-runner/ |
| H A D | mlir-spirv-cpu-runner.cpp | 100 mlir::gpu::GPUDialect, mlir::spirv::SPIRVDialect, in main()
|
| /llvm-project-15.0.7/mlir/lib/Dialect/NVGPU/Transforms/ |
| H A D | OptimizeSharedMemory.cpp | 177 gpu::GPUDialect::getWorkgroupAddressSpace()) in optimizeSharedMemoryReadsAndWrites()
257 gpu::GPUDialect::getWorkgroupAddressSpace()) in runOnOperation()
|
| /llvm-project-15.0.7/mlir/lib/Conversion/GPUToNVVM/ |
| H A D | LowerGpuOpsToNVVMOps.cpp | 189 gpu::GPUDialect::getPrivateAddressSpace()) in runOnOperation()
225 target.addIllegalDialect<gpu::GPUDialect>(); in configureGpuToNVVMConversionLegality()
|
| /llvm-project-15.0.7/mlir/lib/Dialect/GPU/ |
| H A D | CMakeLists.txt | 23 IR/GPUDialect.cpp
|
| /llvm-project-15.0.7/mlir/lib/Conversion/GPUToSPIRV/ |
| H A D | GPUToSPIRV.cpp | 265 if (!gpu::GPUDialect::isKernel(funcOp)) in matchAndRewrite()
296 rewriter.getStringAttr(gpu::GPUDialect::getKernelFuncAttrName())); in matchAndRewrite()
|
| /llvm-project-15.0.7/mlir/lib/Conversion/GPUToROCDL/ |
| H A D | LowerGpuOpsToROCDLOps.cpp | 115 target.addIllegalDialect<gpu::GPUDialect>(); in configureGpuToROCDLConversionLegality()
|
| /llvm-project-15.0.7/utils/bazel/llvm-project-overlay/mlir/ |
| H A D | BUILD.bazel | 510 ":GPUDialect",
2176 ":GPUDialect",
2201 ":GPUDialect",
3617 ":GPUDialect",
3674 ":GPUDialect",
3718 ":GPUDialect",
3769 ":GPUDialect",
3835 ":GPUDialect",
3864 ":GPUDialect",
3890 ":GPUDialect",
[all …]
|
| /llvm-project-15.0.7/mlir/include/mlir/Conversion/ |
| H A D | Passes.td | 677 let dependentDialects = ["gpu::GPUDialect"];
689 let dependentDialects = ["AffineDialect", "gpu::GPUDialect"];
862 "memref::MemRefDialect", "gpu::GPUDialect", "AffineDialect",
|
| /llvm-project-15.0.7/mlir/lib/Conversion/GPUCommon/ |
| H A D | GPUOpsLowering.cpp | 39 /*alignment=*/0, gpu::GPUDialect::getWorkgroupAddressSpace()); in matchAndRewrite()
|
| H A D | GPUToLLVMConversion.cpp | 371 target.addIllegalDialect<gpu::GPUDialect>(); in runOnOperation()
|
| /llvm-project-15.0.7/mlir/test/lib/Dialect/Linalg/ |
| H A D | TestLinalgTransforms.cpp | 50 gpu::GPUDialect>(); in getDependentDialects()
|
| /llvm-project-15.0.7/mlir/test/lib/Dialect/Vector/ |
| H A D | TestVectorTransforms.cpp | 826 registry.insert<scf::SCFDialect, memref::MemRefDialect, gpu::GPUDialect, in MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID()
|
| /llvm-project-15.0.7/mlir/include/mlir/Dialect/GPU/IR/ |
| H A D | GPUOps.td | 261 GPUDialect::getKernelFuncAttrName()) != nullptr;
456 friend LogicalResult GPUDialect::verifyOperationAttribute(Operation *,
|