Home
last modified time | relevance | path

Searched refs:launch_op (Results 1 – 8 of 8) sorted by relevance

/external/tensorflow/tensorflow/compiler/mlir/tools/kernel_gen/transforms/
Dtf_kernel_to_llvm_pass.cc58 Value generateParamsArray(gpu::LaunchFuncOp launch_op,
64 gpu::LaunchFuncOp launch_op, ArrayRef<Value> operands,
97 gpu::LaunchFuncOp launch_op, ArrayRef<Value> operands, in generateParamsArray() argument
99 auto loc = launch_op.getLoc(); in generateParamsArray()
100 auto num_kernel_operands = launch_op.getNumKernelOperands(); in generateParamsArray()
102 loc, launch_op.getOperands().take_back(num_kernel_operands), in generateParamsArray()
145 gpu::LaunchFuncOp launch_op, ArrayRef<Value> operands, in matchAndRewrite() argument
147 if (!launch_op.asyncDependencies().empty() || launch_op.asyncToken()) { in matchAndRewrite()
149 launch_op, "Cannot convert with async dependency or result."); in matchAndRewrite()
152 Location loc = launch_op.getLoc(); in matchAndRewrite()
[all …]
/external/tensorflow/tensorflow/compiler/mlir/tensorflow/transforms/
Doutside_compiled_to_host_launch.cc64 auto launch_op = builder.create<tf_device::LaunchOp>( in WrapOpInLaunch() local
67 host_op->replaceAllUsesWith(launch_op); in WrapOpInLaunch()
69 launch_op.body().push_back(new Block); in WrapOpInLaunch()
70 builder.setInsertionPointToEnd(&launch_op.GetBody()); in WrapOpInLaunch()
75 MLIRContext* context = launch_op.getContext(); in WrapOpInLaunch()
Dcluster_formation.cc104 tf_device::LaunchOp launch_op) { in ReplaceLiveOutExternalUses() argument
105 Region* launch_op_region = &launch_op.body(); in ReplaceLiveOutExternalUses()
106 for (const auto& p : llvm::zip(live_outs, launch_op.getResults())) { in ReplaceLiveOutExternalUses()
170 tf_device::LaunchOp launch_op = builder->create<tf_device::LaunchOp>( in BuildLaunchForCluster() local
175 launch_op.body().takeBody(region); in BuildLaunchForCluster()
179 ReplaceLiveOutExternalUses(live_outs, launch_op); in BuildLaunchForCluster()
Dlaunch_to_device_attribute.cc111 Operation* launch_op = launch.getOperation(); in HoistOpsAndAnnotateWithDevice() local
112 launch_op->getBlock()->getOperations().splice( in HoistOpsAndAnnotateWithDevice()
113 launch_op->getIterator(), launch.GetBody().getOperations(), body.begin(), in HoistOpsAndAnnotateWithDevice()
Dtpu_extract_outside_compilation.cc186 auto launch_op = builder.create<tf_device::LaunchOp>( in CreateLaunchOpForOutsideCluster() local
190 launch_op.body().push_back(new Block); in CreateLaunchOpForOutsideCluster()
191 builder.setInsertionPointToEnd(&launch_op.GetBody()); in CreateLaunchOpForOutsideCluster()
195 return launch_op; in CreateLaunchOpForOutsideCluster()
Dtpu_rewrite_pass.cc696 tf_device::LaunchOp launch_op = AssignDevicesToReplicatedExecute( in Rewrite() local
698 cluster_func.replaceAllUsesWith(launch_op); in Rewrite()
Dshape_inference.cc945 if (auto launch_op = dyn_cast<tf_device::LaunchOp>(op)) { in InferShapeForNonTFDialectOperation() local
946 auto terminator = launch_op.GetBody().getTerminator(); in InferShapeForNonTFDialectOperation()
/external/tensorflow/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/
Dlhlo_legalize_to_gpu.cc90 auto launch_op = rewriter.create<mlir::gpu::LaunchOp>( in matchAndRewrite() local
94 rewriter.setInsertionPointToEnd(&launch_op.body().front()); in matchAndRewrite()
95 auto index = launch_op.getThreadIds().x; in matchAndRewrite()
127 OpFoldResult offset = launch_op.getThreadIds().x; in matchAndRewrite()
142 : launch_op.getThreadIds().x; in matchAndRewrite()
163 rewriter.setInsertionPointToEnd(&launch_op.body().front()); in matchAndRewrite()