tensorflow/compiler/mlir/xla/transforms/legalize_tf_communication.cc
// Rewrites a `mhlo.while` op to receive and forward a `mhlo.token`. Operands to
// the op for all of its regions are extended to have an extra operand `token`.
Rewrites a `mhlo.while` op to receive and forward a `mhlo.token`. Operands to
the op for all of its regions are extended to have an extra operand `token`.
Rewrites a ` mhlo . while ` op to receive and forward a ` mhlo . token ` . Operands to the op for all of its regions are extended to have an extra operand ` token ` .
void RewriteRegionWhileOp(OpBuilder& builder, WhileOp region_while,
                          SmallVectorImpl<OpVisitorState>& ops_to_visit,
                          Value token) {
  llvm::SmallDenseMap<Value, Value> rewritten_operands;
  auto new_val_operands =
      GetValueWithToken(builder, llvm::to_vector(region_while.getOperands()),
                        token, rewritten_operands,
                        true);
  auto new_result_types =
      GetTypeWithToken(builder, llvm::to_vector(region_while.getResultTypes()),
                        true);
  auto new_while = builder.create<WhileOp>(region_while.getLoc(),
                                           new_result_types, new_val_operands);
  new_while.cond().takeBody(region_while.cond());
  new_while.body().takeBody(region_while.body());
  SmallVector<Value> old_while_results = region_while.getResults();
  SmallVector<Value> new_while_results = new_while.getResults();
  ReplaceWithTupleResult(builder, old_while_results, new_while_results,
                          true);
  auto new_token = new_while_results[new_while_results.size() - 1];
  region_while.erase();
  ops_to_visit.push_back({0, new_token, new_while});
}
void RewriteRegionWhileOp ( OpBuilder & builder , WhileOp region_while , SmallVectorImpl < OpVisitorState > & ops_to_visit , Value token ) { llvm :: SmallDenseMap < Value , Value > rewritten_operands ; auto new_val_operands = GetValueWithToken ( builder , llvm :: to_vector ( region_while . getOperands ( ) ) , token , rewritten_operands , true ) ; auto new_result_types = GetTypeWithToken ( builder , llvm :: to_vector ( region_while . getResultTypes ( ) ) , true ) ; auto new_while = builder . create < WhileOp > ( region_while . getLoc ( ) , new_result_types , new_val_operands ) ; new_while . cond ( ) . takeBody ( region_while . cond ( ) ) ; new_while . body ( ) . takeBody ( region_while . body ( ) ) ; SmallVector < Value > old_while_results = region_while . getResults ( ) ; SmallVector < Value > new_while_results = new_while . getResults ( ) ; ReplaceWithTupleResult ( builder , old_while_results , new_while_results , true ) ; auto new_token = new_while_results [ new_while_results . size ( ) - 1 ] ; region_while . erase ( ) ; ops_to_visit . push_back ( { 0 , new_token , new_while } ) ; }
Rewrites a `mhlo.while` op to receive and forward a `mhlo.token`.
Rewrites a ` mhlo . while ` op to receive and forward a ` mhlo . token ` .
// Rewrite region operand to have an extra operand `token`. /*flatten_tuple=*/ /*flatten_tuple*/ // Create new `mhlo.while` op with extra token operand and result. // Move all regions from the old `mhlo.while` op to its replacement. // Forward result from old `mhlo.while` with replacement. /*flatten_tuple*/ // Next op to visit. The replacement is visited but at its first region. The // token result of the new region if is propagated. /*region_idx=*/