JMP	asancall<>(SB)

// Switches SP to g0 stack and calls (AX). Arguments already set.
TEXT	asancall<>(SB), NOSPLIT, $0-0
	get_tls(R12)
	MOVQ	g(R12), R14
	MOVQ	SP, R12		// callee-saved, preserved across the CALL
	CMPQ	R14, $0
	JE	call	// no g; still on a system stack

	MOVQ	g_m(R14), R13
	// Switch to g0 stack.
	MOVQ	m_g0(R13), R10
	CMPQ	R10, R14
	JE	call	// already on g0

	MOVQ	(g_sched+gobuf_sp)(R10), SP
call:

// RUN: odml-to-stablehlo-opt %s -stablehlo-custom-call-legalize-composite | FileCheck %s

// CHECK-LABEL: module
module {
  // CHECK-LABEL: @main
  func.func @main(%arg0: tensor<1xf32>, %arg1: tensor<2xf32>) {
    // CHECK: stablehlo.custom_call @foo
    stablehlo.custom_call @foo() : () -> ()
    // CHECK-NOT: stablehlo.custom_call

// It will first check if the callee was previously handled, and try to reuse
// that result if so. Otherwise, it will clone and convert the callee function,
// and performs stack ops decomposition on it.
template <typename CallOp>
LogicalResult HandlePartitionedCallOp(
    CallOp call, func::FuncOp callee, ModuleOp module,
    const llvm::SmallDenseMap<Value, Value>& data_var_to_size_var,

	MOVL	CX, 12(SP)
	MOVL	$0, 16(SP)
	CALL	runtime·reflectcallmove(SB)
	RET

CALLFN(·call16, 16)
CALLFN(·call32, 32)
CALLFN(·call64, 64)
CALLFN(·call128, 128)
CALLFN(·call256, 256)
CALLFN(·call512, 512)
CALLFN(·call1024, 1024)
CALLFN(·call2048, 2048)
CALLFN(·call4096, 4096)
CALLFN(·call8192, 8192)
CALLFN(·call16384, 16384)
CALLFN(·call32768, 32768)
CALLFN(·call65536, 65536)
CALLFN(·call131072, 131072)

  // CHECK: "tf.StatefulPartitionedCall"
  // CHECK-SAME: f = @callee
  %call = "tf.StatefulPartitionedCall"(%stack, %arg0) {f = @callee, config = "", config_proto = "", executor_type = ""}
    : (tensor<!tf_type.resource>, tensor<i1>) -> tensor<!tf_type.resource>
  // CHECK: "tf.PartitionedCall"
  // CHECK-SAME: f = @callee
  %call2 = "tf.PartitionedCall"(%stack, %arg0) {f = @callee, config = "", config_proto = "", executor_type = ""}

// func runtime·RaceRead(addr uintptr)
TEXT	runtime·RaceRead(SB), NOSPLIT, $0-8
	// This needs to be a tail call, because raceread reads caller pc.
	JMP	runtime·raceread(SB)

// func runtime·racereadpc(void *addr, void *callpc, void *pc)
TEXT	runtime·racereadpc(SB), NOSPLIT, $0-24
	// void __tsan_read_pc(ThreadState *thr, void *addr, void *callpc, void *pc);
	MOVD	$__tsan_read_pc(SB), R1
	LMG	addr+0(FP), R3, R5
	JMP	racecalladdr<>(SB)

	// Allocate space for argument storage to call runtime.sigtrampgo.
	STACK_AND_SAVE_HOST_TO_GO_ABI(32)

	// this might be called in external code context,
	// where g is not set.
	MOVBZ	runtime·iscgo(SB), R6
	CMP	R6, $0
	BEQ	2(PC)
	BL	runtime·load_g(SB)

	// R3,R4,R5 already hold the arguments. Forward them on.

  if (lowered_callee != callee) {
    if (!info.signature_change) {
      // Signature is not modified. We do not need to keep two copies.
      lowered_callee.setName(
          StringAttr::get(callee->getContext(), callee.getName()));
      callee.erase();
    } else {
      // Add the clone with a new name.
      lowered_callee.setName(StringAttr::get(
          callee->getContext(),

  }

  // Clone the callee before making changes.
  SmallString<64> name_base = callee.getName();
  auto module = callee->getParentOfType<ModuleOp>();
  name_base += "_resource_lifted";
  auto name = name_base;
  callee = callee.clone();
  callee.setPrivate();
  callee.setName(mlir::StringAttr::get(callee->getContext(), name));
  SymbolTable(module).insert(callee);
  result->lifted_callee = callee;

    def "one arg closure called with subject and no inputs"() {
        given:
        def called = false
        String thing = "1"
        def closure = { String val ->
            called = true
            assert val.is(thing)
            assert delegate.is(thing)
        }

        when:
        action(closure).execute(thing, [])

        then:
        called
    }