| /external/tensorflow/tensorflow/compiler/mlir/tools/kernel_gen/transforms/ |
| D | kernel_lowering_passes.cc | 32 using gpu::GPUModuleOp;
49 GPUModuleOp m = getOperation(); in runOnOperation()
76 gpu::GPUModuleOp m = getOperation(); in runOnOperation()
93 std::unique_ptr<OperationPass<GPUModuleOp> > CreateGpuKernelToNvvmPass() { in CreateGpuKernelToNvvmPass()
97 std::unique_ptr<OperationPass<GPUModuleOp> > CreateGpuKernelToRocdlPass() { in CreateGpuKernelToRocdlPass()
|
| D | passes.h | 69 std::unique_ptr<OperationPass<gpu::GPUModuleOp>> CreateGpuKernelToBlobPass(
92 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>>
96 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>>
|
| D | tf_kernel_to_llvm_pass.cc | 156 auto kernel_module = SymbolTable::lookupNearestSymbolFrom<gpu::GPUModuleOp>( in matchAndRewrite()
272 target.addLegalOp<ModuleOp, ModuleTerminatorOp, gpu::GPUModuleOp>(); in runOnOperation()
274 target.markOpRecursivelyLegal<gpu::GPUModuleOp>(); in runOnOperation()
281 for (auto op : llvm::make_early_inc_range(m.getOps<gpu::GPUModuleOp>())) { in runOnOperation()
|
| D | gpu_kernel_to_blob_pass.cc | 67 mlir::gpu::GPUModuleOp gpu_module = getOperation(); in runOnOperation()
82 mlir::gpu::GPUModuleOp gpu_module) { in GetGpuBinaryBlob()
254 std::unique_ptr<OperationPass<gpu::GPUModuleOp>> CreateGpuKernelToBlobPass( in CreateGpuKernelToBlobPass()
|
| D | passes.td | 64 def GpuKernelToNVVMPass : Pass<"gpu-kernel-to-nvvm", "gpu::GPUModuleOp"> {
69 def GpuKernelToROCDLPass : Pass<"gpu-kernel-to-rocdl", "gpu::GPUModuleOp"> {
74 def GpuKernelToBlobPass : Pass<"gpu-kernel-to-blob", "gpu::GPUModuleOp"> {
|
| /external/tensorflow/tensorflow/compiler/xla/service/mlir_gpu/ |
| D | kernel_lowering.cc | 152 LowerToNVVMPass, ::mlir::OperationPass<::mlir::gpu::GPUModuleOp>> {
159 ::mlir::gpu::GPUModuleOp m = getOperation(); in runOnOperation()
185 auto& kernelPm = pm.nest<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToNVVM()
206 LowerToROCDLPass, ::mlir::OperationPass<::mlir::gpu::GPUModuleOp>> {
213 ::mlir::gpu::GPUModuleOp m = getOperation(); in runOnOperation()
255 auto& kernelPm = pm.nest<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToROCDL()
274 module.walk([&kernelModule](mlir::gpu::GPUModuleOp nestedModule) { in ExtractKernelModule()
|
| /external/llvm-project/mlir/include/mlir/Conversion/GPUToROCDL/ |
| D | GPUToROCDLPass.h | 22 class GPUModuleOp; variable
32 std::unique_ptr<OperationPass<gpu::GPUModuleOp>>
|
| /external/llvm-project/mlir/include/mlir/Conversion/GPUToNVVM/ |
| D | GPUToNVVMPass.h | 22 class GPUModuleOp; variable
35 std::unique_ptr<OperationPass<gpu::GPUModuleOp>> createLowerGpuOpsToNVVMOpsPass(
|
| /external/llvm-project/mlir/lib/Conversion/GPUToROCDL/ |
| D | LowerGpuOpsToROCDLOps.cpp | 53 gpu::GPUModuleOp m = getOperation(); in runOnOperation()
79 target.addLegalOp<gpu::YieldOp, gpu::GPUModuleOp, gpu::ModuleEndOp>(); in runOnOperation()
122 std::unique_ptr<OperationPass<gpu::GPUModuleOp>>
|
| /external/llvm-project/mlir/test/lib/Dialect/SPIRV/ |
| D | TestEntryPointAbi.cpp | 25 OperationPass<gpu::GPUModuleOp>> {
43 gpu::GPUModuleOp gpuModule = getOperation(); in runOnOperation()
|
| /external/llvm-project/mlir/lib/Conversion/GPUCommon/ |
| D | ConvertKernelFuncToBlob.cpp | 48 : public PassWrapper<GpuKernelToBlobPass, OperationPass<gpu::GPUModuleOp>> {
59 gpu::GPUModuleOp module = getOperation(); in runOnOperation()
154 std::unique_ptr<OperationPass<gpu::GPUModuleOp>>
|
| D | ConvertLaunchFuncToRuntimeCalls.cpp | 260 class EraseGpuModuleOpPattern : public OpRewritePattern<gpu::GPUModuleOp> {
261 using OpRewritePattern<gpu::GPUModuleOp>::OpRewritePattern;
263 LogicalResult matchAndRewrite(gpu::GPUModuleOp op, in matchAndRewrite()
586 auto kernelModule = SymbolTable::lookupNearestSymbolFrom<gpu::GPUModuleOp>( in matchAndRewrite()
|
| /external/llvm-project/mlir/include/mlir/Conversion/GPUCommon/ |
| D | GPUCommonPass.h | 28 class GPUModuleOp; variable
77 std::unique_ptr<OperationPass<gpu::GPUModuleOp>>
|
| /external/llvm-project/mlir/lib/Dialect/GPU/Transforms/ |
| D | KernelOutlining.cpp | 278 gpu::GPUModuleOp createKernelModule(gpu::GPUFuncOp kernelFunc, in createKernelModule()
287 gpu::GPUModuleOp::getOperationName()); in createKernelModule()
288 gpu::GPUModuleOp::build(builder, state, kernelFunc.getName()); in createKernelModule()
289 auto kernelModule = cast<gpu::GPUModuleOp>(Operation::create(state)); in createKernelModule()
|
| /external/llvm-project/mlir/lib/Conversion/GPUToNVVM/ |
| D | LowerGpuOpsToNVVMOps.cpp | 108 gpu::GPUModuleOp m = getOperation(); in runOnOperation()
156 target.addLegalOp<gpu::YieldOp, gpu::GPUModuleOp, gpu::ModuleEndOp>(); in configureGpuToNVVMConversionLegality()
202 std::unique_ptr<OperationPass<gpu::GPUModuleOp>>
|
| /external/llvm-project/mlir/lib/Conversion/GPUToSPIRV/ |
| D | ConvertGPUToSPIRV.cpp | 76 class GPUModuleConversion final : public SPIRVOpLowering<gpu::GPUModuleOp> {
78 using SPIRVOpLowering<gpu::GPUModuleOp>::SPIRVOpLowering;
81 matchAndRewrite(gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
281 gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands, in matchAndRewrite()
|
| D | ConvertGPUToSPIRVPass.cpp | 48 module.walk([&builder, &kernelModules](gpu::GPUModuleOp moduleOp) { in runOnOperation()
|
| /external/llvm-project/mlir/lib/Conversion/ |
| D | PassDetail.h | 24 class GPUModuleOp; variable
|
| /external/tensorflow/tensorflow/compiler/mlir/tools/kernel_gen/ |
| D | kernel_creator.cc | 344 auto gpu_modules = module.getOps<mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToLowLevelIr()
361 auto& kernelPm = pm.nest<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToLowLevelIr()
401 auto& kernel_pm = pm.nest<mlir::gpu::GPUModuleOp>(); in GenerateDeviceCode()
|
| /external/llvm-project/mlir/lib/Dialect/GPU/IR/ |
| D | GPUDialect.cpp | 95 auto kernelModule = module.lookupSymbol<GPUModuleOp>(kernelModuleName); in verifyOperationAttribute()
437 auto kernelModule = kernelFunc->getParentOfType<GPUModuleOp>(); in build()
807 void GPUModuleOp::build(OpBuilder &builder, OperationState &result, in build()
831 GPUModuleOp::ensureTerminator(*body, parser.getBuilder(), result.location); in parseGPUModuleOp()
835 static void print(OpAsmPrinter &p, GPUModuleOp op) { in print()
|
| /external/llvm-project/mlir/tools/mlir-cuda-runner/ |
| D | mlir-cuda-runner.cpp | 114 auto &kernelPm = pm.nest<gpu::GPUModuleOp>(); in runMLIRPasses()
|
| /external/llvm-project/mlir/lib/Conversion/GPUToVulkan/ |
| D | ConvertGPULaunchFuncToVulkanLaunchFunc.cpp | 92 llvm::make_early_inc_range(getOperation().getOps<gpu::GPUModuleOp>())) in runOnOperation()
|
| /external/llvm-project/mlir/lib/Conversion/SPIRVToLLVM/ |
| D | ConvertLaunchFuncToLLVMCalls.cpp | 272 llvm::make_early_inc_range(module.getOps<gpu::GPUModuleOp>())) in runOnOperation()
|
| /external/llvm-project/mlir/tools/mlir-rocm-runner/ |
| D | mlir-rocm-runner.cpp | 311 auto &kernelPm = pm.nest<gpu::GPUModuleOp>(); in runMLIRPasses()
|
| /external/llvm-project/mlir/include/mlir/Conversion/ |
| D | Passes.td | 115 def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {
130 def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
|