Searched refs:GPUModuleOp (Results 1 – 15 of 15) sorted by relevance
37 using ::mlir::gpu::GPUModuleOp;65 pm.addNestedPass<GPUModuleOp>(createForLoopSpecializationPass()); in createHloToGpuPipeline()66 pm.addNestedPass<GPUModuleOp>(createLowerAffinePass()); in createHloToGpuPipeline()67 pm.addNestedPass<GPUModuleOp>(createCanonicalizerPass()); in createHloToGpuPipeline()68 pm.addNestedPass<GPUModuleOp>(createConvertSCFToCFPass()); in createHloToGpuPipeline()71 pm.addNestedPass<GPUModuleOp>(createGpuKernelToRocdlPass()); in createHloToGpuPipeline()73 pm.addNestedPass<GPUModuleOp>(createGpuKernelToNvvmPass()); in createHloToGpuPipeline()78 pm.addNestedPass<GPUModuleOp>(createStripDebugInfoPass()); in createHloToGpuPipeline()
91 std::unique_ptr<OperationPass<gpu::GPUModuleOp> >96 std::unique_ptr<OperationPass<gpu::GPUModuleOp> >
223 auto callback = [&](gpu::GPUModuleOp gpuModule) -> WalkResult { in runOnOperation()
246 for (auto gpuModuleOp : moduleOp.getBodyRegion().getOps<gpu::GPUModuleOp>()) { in matchAndRewrite()
27 class GPUModuleOp; variable40 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>>44 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>>
26 def GpuKernelToNVVMPass : Pass<"gpu-kernel-to-nvvm", "gpu::GPUModuleOp"> {32 def GpuKernelToROCDLPass : Pass<"gpu-kernel-to-rocdl", "gpu::GPUModuleOp"> {
27 class GPUModuleOp; variable
165 auto kernel_module = SymbolTable::lookupNearestSymbolFrom<gpu::GPUModuleOp>( in matchAndRewrite()289 target.addLegalOp<ModuleOp, gpu::GPUModuleOp>(); in runOnOperation()291 target.markOpRecursivelyLegal<gpu::GPUModuleOp>(); in runOnOperation()300 for (auto op : llvm::make_early_inc_range(m.getOps<gpu::GPUModuleOp>())) { in runOnOperation()
62 gpu::GPUModuleOp gpu_module = getOperation(); in runOnOperation()77 gpu::GPUModuleOp gpu_module) { in GetGpuBinaryBlob()259 std::unique_ptr<OperationPass<gpu::GPUModuleOp>> CreateGpuKernelToBlobPass( in CreateGpuKernelToBlobPass()
70 std::unique_ptr<OperationPass<gpu::GPUModuleOp>> CreateGpuKernelToBlobPass(
80 def GpuKernelToBlobPass : Pass<"gpu-kernel-to-blob", "gpu::GPUModuleOp"> {
317 auto gpu_modules = module.getOps<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToLowLevelIr()318 for (::mlir::gpu::GPUModuleOp gpu_module : gpu_modules) { in LowerKernelBodiesToLowLevelIr()334 auto& kernelPm = pm.nest<::mlir::gpu::GPUModuleOp>(); in LowerKernelBodiesToLowLevelIr()378 auto& kernel_pm = pm.nest<mlir::gpu::GPUModuleOp>(); in GenerateDeviceCode()
82 using mlir::gpu::GPUModuleOp;154 class GpuModuleOpLowering : public OpRewritePattern<GPUModuleOp> {158 LogicalResult matchAndRewrite(GPUModuleOp op, in matchAndRewrite()373 auto gpu_module = kernel->getParentOfType<mlir::gpu::GPUModuleOp>(); in matchAndRewrite()
341 auto gpu_module = rewriter.create<mlir::gpu::GPUModuleOp>(loc, "gpu_module"); in Rewrite()
1600 auto gpu_module = kernel_func->getParentOfType<mlir::gpu::GPUModuleOp>(); in EmitLaunchFunc()