OptionalAttr<DictArrayAttr>:$res_attrs
  );

  let results = (outs);

  // When the regions is empty, the tfr.func is an external function and used
  // to model the element type constraints of the tf op. Otherwise, there is one
  // region containing the composition.
  let regions = (region VariadicRegion<AnyRegion>:$body);

  let skipDefaultBuilders = 1;

  let builders = [

        }
        // Transfer the annotation to both region arguments
        for (Region* region : while_region.getRegions()) {
          BlockArgument arg = region->getArgument(output_index);
          LLVM_DEBUG(llvm::dbgs()
                     << "  Propagating device = '" << *device << "' to arg #"
                     << output_index << " of region #"
                     << region->getRegionNumber() << "\n");

			return ""
		}
		old, new := st.Goroutine()
		return fmt.Sprintf("%s -> %s", old, new)
	case trace.EventRegionBegin:
		return fmt.Sprintf("region %q begin", ev.Region().Type)
	case trace.EventRegionEnd:
		return fmt.Sprintf("region %q end", ev.Region().Type)
	case trace.EventTaskBegin:
		t := ev.Task()
		return fmt.Sprintf("task %q (D %d, parent %d) begin", t.Type, t.ID, t.Parent)
	case trace.EventTaskEnd:

		return shadowRange(lo + hi<<16)
	}
	if hi < sr.lo() || lo > sr.hi() {
		// The two regions don't overlap or abut, so we would
		// have to keep track of multiple disjoint ranges.
		// Because we can only keep one, keep the larger one.
		if sr.hi()-sr.lo() >= hi-lo {
			return sr
		}
		return shadowRange(lo + hi<<16)
	}
	// Regions overlap or abut - compute the union.
	return shadowRange(min(lo, sr.lo()) + max(hi, sr.hi())<<16)

      dependencies.push_back(dependency->getParentOp());
    if (!safe_to_merge) break;
  }
  return safe_to_merge;
}

// Move the body excluding the terminators of else and then regions from
// 'second_if' to 'first_if'.
void MoveBranches(TF::IfRegionOp first_if, TF::IfRegionOp second_if) {
  Block& first_if_then_block = first_if.getThenBranch().front();

func (r Region) String() string {
	if r < isoRegionOffset {
		if r == 0 {
			return "ZZ"
		}
		return fmt.Sprintf("%03d", r.M49())
	}
	r -= isoRegionOffset
	return regionISO.Elem(int(r))[:2]
}

// ISO3 returns the 3-letter ISO code of r.
// Note that not all regions have a 3-letter ISO code.
// In such cases this method returns "ZZZ".
func (r Region) ISO3() string {

  auto& region = op->getRegion(region_idx);
  region.getArgument(0).replaceAllUsesWith(implicit_operand);
  region.front().eraseArguments(0, region.getNumArguments());
}

// Rewrites an `mhlo.if` op or its region. If `region_idx` is not set, the op
// operands and results are rewritten. If `region_idx` is set, region

			mode := ev.Args[1]
			regions := activeRegions[ev.G]
			if mode == 0 { // region start
				activeRegions[ev.G] = append(regions, ev) // push
			} else if mode == 1 { // region end
				n := len(regions)
				if n > 0 { // matching region start event is in the trace.
					s := regions[n-1]
					if s.Args[0] != ev.Args[0] || s.Args[2] != ev.Args[2] { // task id, region name mismatch

  }
}

// An override that takes region.
void PropagateUsage(
    Region* region, ElementsAttr value, int argument_index,
    llvm::SmallVector<std::pair<Region*, int>, 4>* work_list,
    llvm::MapVector<Operation*, llvm::SmallVector<unsigned int, 4>>*
        arguments_to_erase) {
  auto arg = region->getArgument(argument_index);
  for (auto& usage : arg.getUses()) {

      %res = tfrt_fallback_async.executeop key(0) cost(100) device("/CPU:0") "some.op"(%arg) : 1

    Note that the trailing number indicates the number of results.

    The func attrs and op attrs are placed in two separate regions.

    Example:
      %res = tfrt_fallback_async.executeop key(0) device("/CPU:0") "some.op"(%arg)
          {attr1 = value, attr2 = value2} {f_attr1 = "fn1", f_attr2 = "fn2"} : 1