Home
last modified time | relevance | path

Searched refs:LaunchOp (Results 1 – 23 of 23) sorted by relevance

/external/tensorflow/tensorflow/compiler/mlir/tensorflow/transforms/
Dtpu_dynamic_layout_pass.cc143 tf_device::LaunchOp compile_launch, in BuildGetLayout()
157 TF::TPUCopyWithLayoutOp BuildCopyWithLayout(tf_device::LaunchOp execute_launch, in BuildCopyWithLayout()
158 tf_device::LaunchOp compile_launch, in BuildCopyWithLayout()
168 TF::TPUExecuteOp execute, tf_device::LaunchOp execute_launch, in HandleInput()
169 tf_device::LaunchOp compile_launch) { in HandleInput()
184 tf_device::LaunchOp execute_launch, tf_device::LaunchOp compile_launch, in HandleReplicatedInputs()
221 tf_device::LaunchOp compile_launch, in HandleCompileAndExecutes()
222 llvm::MutableArrayRef<tf_device::LaunchOp> execute_launches, in HandleCompileAndExecutes()
286 llvm::dyn_cast<tf_device::LaunchOp>(compile->getParentOp()); in runOnFunction()
289 llvm::SmallVector<tf_device::LaunchOp, 4> execute_launches; in runOnFunction()
[all …]
Dtpu_merge_variables_with_execute.cc126 tf_device::LaunchOp execute_launch, bool check_device, in BuildVariableAccessInfo()
332 tf_device::LaunchOp execute_launch, in ReplaceParallelExecute()
333 tf_device::LaunchOp merged_execute_launch, in ReplaceParallelExecute()
407 void ReplaceExecute(tf_device::LaunchOp execute_launch, in ReplaceExecute()
408 tf_device::LaunchOp merged_execute_launch, in ReplaceExecute()
424 TF::_TPUCompileMlirOp GetTPUCompileOp(tf_device::LaunchOp execute_launch) { in GetTPUCompileOp()
428 auto compile_launch = llvm::dyn_cast_or_null<tf_device::LaunchOp>( in GetTPUCompileOp()
437 LogicalResult UpdateSerializedModule(tf_device::LaunchOp execute_launch, in UpdateSerializedModule()
486 LogicalResult MergeForOneTPUExecute(tf_device::LaunchOp execute_launch, in MergeForOneTPUExecute()
552 auto merged_execute_launch = builder->create<tf_device::LaunchOp>( in MergeForOneTPUExecute()
[all …]
Dlaunch_to_device_attribute.cc67 tf_device::LaunchOp launch, in AssignDevicesInRegion()
101 tf_device::LaunchOp launch) { in HoistOpsAndAnnotateWithDevice()
128 auto result = getOperation().walk([&tf_dialect](tf_device::LaunchOp launch) { in runOnFunction()
Dtpu_colocate_composite_resource_ops.cc44 auto launch = builder->create<tf_device::LaunchOp>( in WrapOpInLaunch()
109 llvm::SmallVector<tf_device::LaunchOp, 8> execute_launches; in runOnFunction()
110 getFunction().walk([&](tf_device::LaunchOp op) { in runOnFunction()
Dcluster_formation.cc104 tf_device::LaunchOp launch_op) { in ReplaceLiveOutExternalUses()
170 tf_device::LaunchOp launch_op = builder->create<tf_device::LaunchOp>( in BuildLaunchForCluster()
Dtpu_variable_runtime_reformatting.cc148 tf_device::LaunchOp compile_launch) { in AnnotateCompileOpAndGetExecuteArgToWhileArgsMapping()
368 auto launch = builder->create<tf_device::LaunchOp>( in WrapOpInLaunch()
386 tf_device::LaunchOp execute_launch; in HandleReplicateOp()
388 replicate.GetBody().getOps<tf_device::LaunchOp>()) { in HandleReplicateOp()
408 auto compile_launch = llvm::dyn_cast<tf_device::LaunchOp>(compile); in HandleReplicateOp()
Dtpu_extract_head_tail_outside_compilation.cc86 tf_device::LaunchOp CreateLaunchForBlock(OpBuilder* builder, Operation* op, in CreateLaunchForBlock()
108 auto launch = builder->create<tf_device::LaunchOp>( in CreateLaunchForBlock()
190 tf_device::LaunchOp launch = CreateLaunchForBlock( in CreateHeadComputation()
298 tf_device::LaunchOp launch = CreateLaunchForBlock( in CreateTailComputation()
Dtpu_extract_outside_compilation.cc182 tf_device::LaunchOp CreateLaunchOpForOutsideCluster( in CreateLaunchOpForOutsideCluster()
186 auto launch_op = builder.create<tf_device::LaunchOp>( in CreateLaunchOpForOutsideCluster()
410 void RemoveOutsideCompilation(tf_device::LaunchOp host_launch_op) { in RemoveOutsideCompilation()
436 tf_device::LaunchOp host_launch_op = in CreateParallelExecuteForOutsideCompilation()
Doutside_compiled_to_host_launch.cc64 auto launch_op = builder.create<tf_device::LaunchOp>( in WrapOpInLaunch()
Dtpu_rewrite_pass.cc328 tf_device::LaunchOp WrapOpInLaunch(OpBuilder* builder, Location loc, in WrapOpInLaunch()
332 auto launch = builder->create<tf_device::LaunchOp>( in WrapOpInLaunch()
543 tf_device::LaunchOp AssignDevicesToReplicatedExecute( in AssignDevicesToReplicatedExecute()
696 tf_device::LaunchOp launch_op = AssignDevicesToReplicatedExecute( in Rewrite()
Dreplicate_to_island.cc122 if (auto launch = dyn_cast<tf_device::LaunchOp>(op)) in UpdateRegionReplicateVariantOps()
Dshape_inference.cc114 return isa<tf_device::ReturnOp, tf_device::ClusterOp, tf_device::LaunchOp, in IsSupportedNonTFOp()
945 if (auto launch_op = dyn_cast<tf_device::LaunchOp>(op)) { in InferShapeForNonTFDialectOperation()
/external/llvm-project/mlir/include/mlir/Dialect/GPU/
DUtils.h24 class LaunchOp; variable
35 gpu::GPUFuncOp outlineKernelFunc(gpu::LaunchOp launchOp, StringRef kernelFnName,
41 LogicalResult sinkOperationsIntoLaunchOp(gpu::LaunchOp launchOp);
DGPUOps.td553 def GPU_TerminatorOp : GPU_Op<"terminator", [HasParent<"LaunchOp">,
/external/llvm-project/mlir/lib/Dialect/GPU/IR/
DGPUDialect.cpp245 void LaunchOp::build(OpBuilder &builder, OperationState &result, in build()
262 KernelDim3 LaunchOp::getBlockIds() { in getBlockIds()
268 KernelDim3 LaunchOp::getThreadIds() { in getThreadIds()
274 KernelDim3 LaunchOp::getGridSize() { in getGridSize()
280 KernelDim3 LaunchOp::getBlockSize() { in getBlockSize()
286 KernelDim3 LaunchOp::getGridSizeOperandValues() { in getGridSizeOperandValues()
290 KernelDim3 LaunchOp::getBlockSizeOperandValues() { in getBlockSizeOperandValues()
294 static LogicalResult verify(LaunchOp op) { in verify()
300 LaunchOp::kNumConfigOperands + op.getNumOperands()) in verify()
317 .append("in '", LaunchOp::getOperationName(), "' body region"); in verify()
[all …]
/external/llvm-project/mlir/lib/Dialect/GPU/Transforms/
DKernelOutlining.cpp106 LogicalResult mlir::sinkOperationsIntoLaunchOp(gpu::LaunchOp launchOp) { in sinkOperationsIntoLaunchOp()
138 static gpu::GPUFuncOp outlineKernelFuncImpl(gpu::LaunchOp launchOp, in outlineKernelFuncImpl()
197 gpu::GPUFuncOp mlir::outlineKernelFunc(gpu::LaunchOp launchOp, in outlineKernelFunc()
214 static void convertToLaunchFuncOp(gpu::LaunchOp launchOp, in convertToLaunchFuncOp()
243 auto funcWalkResult = func.walk([&](gpu::LaunchOp op) { in runOnOperation()
DAsyncRegionRewriter.cpp56 if (isa<gpu::LaunchOp>(op)) in operator ()()
/external/llvm-project/mlir/lib/Conversion/GPUToVulkan/
DConvertGPULaunchFuncToVulkanLaunchFunc.cpp113 gpu::LaunchOp::kNumConfigOperands, in declareVulkanLaunchFunc()
169 gpu::LaunchOp::kNumConfigOperands, in convertGpuLaunchFunc()
/external/tensorflow/tensorflow/compiler/mlir/tensorflow/ir/
Dtf_device.cc128 bool LaunchOp::WrapsSingleOp() { return BlockWrapsSingleOp(&GetBody()); } in WrapsSingleOp()
683 struct DropEmptyLaunch : public OpRewritePattern<LaunchOp> {
684 using OpRewritePattern<LaunchOp>::OpRewritePattern;
686 LogicalResult matchAndRewrite(LaunchOp op, in matchAndRewrite()
700 void LaunchOp::getCanonicalizationPatterns(OwningRewritePatternList& results, in getCanonicalizationPatterns()
/external/llvm-project/mlir/lib/Conversion/SCFToGPU/
DSCFToGPU.cpp218 auto launchOp = builder.create<gpu::LaunchOp>( in createLaunch()
378 ParallelOp parallelOp, gpu::LaunchOp launchOp, in processParallelLoop()
577 gpu::LaunchOp launchOp = rewriter.create<gpu::LaunchOp>( in matchAndRewrite()
/external/tensorflow/tensorflow/compiler/mlir/tensorflow/analysis/
Dresource_alias_analysis.cc206 } else if (isa<tf_device::LaunchOp, tf_device::ClusterOp>(op)) { in BacktrackValue()
348 } else if (isa<tf_device::LaunchOp, tf_device::ClusterOp>(op)) { in ResourceAliasAnalysisInfo()
/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()
/external/tensorflow/tensorflow/compiler/mlir/tensorflow/g3doc/
Dspace_to_depth.md129 tf\_device::LaunchOp and get the first Convolution and its shape;