Searched refs:kernelPm (Results 1 – 4 of 4) sorted by relevance
185 auto& kernelPm = pm.nest<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToNVVM() local186 kernelPm.addPass(::mlir::createLowerToCFGPass()); in LowerKernelBodiesToNVVM()187 kernelPm.addPass(absl::make_unique<LowerToNVVMPass>()); in LowerKernelBodiesToNVVM()189 kernelPm.addNestedPass<::mlir::FuncOp>(::mlir::createCanonicalizerPass()); in LowerKernelBodiesToNVVM()190 kernelPm.addNestedPass<::mlir::FuncOp>(::mlir::createCSEPass()); in LowerKernelBodiesToNVVM()255 auto& kernelPm = pm.nest<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToROCDL() local256 kernelPm.addPass(::mlir::createLowerToCFGPass()); in LowerKernelBodiesToROCDL()257 kernelPm.addPass(absl::make_unique<LowerToROCDLPass>()); in LowerKernelBodiesToROCDL()260 kernelPm.addNestedPass<::mlir::FuncOp>(::mlir::createCanonicalizerPass()); in LowerKernelBodiesToROCDL()261 kernelPm.addNestedPass<::mlir::FuncOp>(::mlir::createCSEPass()); in LowerKernelBodiesToROCDL()[all …]
114 auto &kernelPm = pm.nest<gpu::GPUModuleOp>(); in runMLIRPasses() local115 kernelPm.addPass(createStripDebugInfoPass()); in runMLIRPasses()116 kernelPm.addPass(createLowerGpuOpsToNVVMOpsPass()); in runMLIRPasses()117 kernelPm.addPass(createConvertGPUKernelToBlobPass( in runMLIRPasses()
311 auto &kernelPm = pm.nest<gpu::GPUModuleOp>(); in runMLIRPasses() local312 kernelPm.addPass(createStripDebugInfoPass()); in runMLIRPasses()313 kernelPm.addPass(createLowerGpuOpsToROCDLOpsPass()); in runMLIRPasses()314 kernelPm.addPass(createConvertGPUKernelToBlobPass( in runMLIRPasses()
361 auto& kernelPm = pm.nest<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToLowLevelIr() local362 kernelPm.addPass(::mlir::createLowerToCFGPass()); in LowerKernelBodiesToLowLevelIr()364 kernelPm.addPass(mlir::kernel_gen::transforms::CreateGpuKernelToRocdlPass()); in LowerKernelBodiesToLowLevelIr()366 kernelPm.addPass(mlir::kernel_gen::transforms::CreateGpuKernelToNvvmPass()); in LowerKernelBodiesToLowLevelIr()