skip--
			} else {
				// Callers expect the pc buffer to contain return addresses
				// and do the -1 themselves, so we add 1 to the call pc to
				// create a "return pc". Since there is no actual call, here
				// "return pc" just means a pc you subtract 1 from to get
				// the pc of the "call". The actual no-op we insert may or
				// may not be 1 byte.
				pcBuf[n] = uf.pc + 1
				n++
			}
			u.calleeFuncID = sf.funcID

	BL.GE	4(PC)  // CALL.GE 4(PC)    // 020000ab
	BL.LT	3(PC)  // CALL.LT 3(PC)    // 010000bb
	BL.GT	2(PC)  // CALL.GT 2(PC)    // 000000cb
	BL.LE	1(PC)  // CALL.LE 1(PC)    // ffffffdb
	ADD	$0, R0, R0
	BL	-1(PC) // CALL -1(PC)      // fdffffeb
	BL	-2(PC) // CALL -2(PC)      // fcffffeb
	BL	-3(PC) // CALL -3(PC)      // fbffffeb
	BL	-4(PC) // CALL -4(PC)      // faffffeb

	XXPERMDI X0, X0, $2, X0
	XXPERMDI X1, X1, $2, X1
	CALL   p256MulInternal<>(SB)

	VOR T0, T0, Z3L
	VOR T1, T1, Z3H

	// X=Y;  Y- ;  MUL; X=T  // T3 = T1*T1         T2
	VOR  Y0, Y0, X0
	VOR  Y1, Y1, X1
	CALL p256MulInternal<>(SB)
	VOR  T0, T0, X0
	VOR  T1, T1, X1

	// X- ;  Y- ;  MUL; T4=T // T4 = T3*T1         T2        T4
	CALL p256MulInternal<>(SB)
	VOR  T0, T0, T4L
	VOR  T1, T1, T4H

	// Reserve PLT entries for these symbols and generate call
	// stubs. The call stubs need to live in .text, which is why we
	// need to do this pass this early.

	// Reserve PLT entry and generate symbol resolver
	addpltsym(ctxt, ldr, r.Sym())

	// The stub types are described in gencallstub.
	stubType := 0
	stubTypeStr := ""

	// For now, the choice of call stub type is determined by whether

   * is accepted, then this future is guaranteed to have been completed with the supplied future by
   * the time this method returns. If the supplied future is not done and the call is accepted, then
   * the future will be <i>set asynchronously</i>. Note that such a result, though not yet known,
   * cannot be overridden by a call to a {@code set*} method, only by a call to {@link #cancel}.
   *

   * is accepted, then this future is guaranteed to have been completed with the supplied future by
   * the time this method returns. If the supplied future is not done and the call is accepted, then
   * the future will be <i>set asynchronously</i>. Note that such a result, though not yet known,
   * cannot be overridden by a call to a {@code set*} method, only by a call to {@link #cancel}.
   *

    // returned by _Retval nodes.
    std::unique_ptr<Graph> graph_;

    // Which device are these nodes on? Used to assign a device to the call
    // node.
    string device_;

    // NodeDef for the function call node.
    NodeDef call_node_def_;

    // Name that is used for the call node. This may not be
    // call_node_def_.name() if the client supplies a rewrite lambda.
    string function_def_name_;

// CHECK: %[[CONST_0:.+]] = stablehlo.constant dense<3.000000e-01> : tensor<2x3xf32>
// CHECK: %[[CALL:.+]] = call @composite_dot_general_fn(%[[ARG_0]], %[[CONST_0]]) : (tensor<1x2xf32>, tensor<2x3xf32>) -> tensor<1x3xf32>
// CHECK: return %[[CALL]]

  func.func private @composite_dot_general_fn(%arg0: tensor<1x2xf32>, %arg1: tensor<2x3xf32>) -> tensor<1x3xf32> attributes {_from_xla_call_module} {

// Populates desiredStateOfWorld cache with two ode/volume/pod tuples pointing to the created volume
// Calls Run()
// Verifies there is one attach call and no detach calls.
// Deletes the node/volume/pod tuple from desiredStateOfWorld which succeeded in attaching
// Verifies there are two attach call and one detach call.
func Test_Run_OneVolumeAttachAndDetachMultipleNodesWithReadWriteOnce(t *testing.T) {
	// Arrange

// CHECK:           %[[VAL_4:.*]]:2 = call @func_0_GPU_FLOAT(%[[VAL_0]], %[[VAL_1]], %[[VAL_2]], %[[VAL_3]]) {tac.device = "GPU", tac.inference_type = "FLOAT", tac.interface_name = "func_0"} : (tensor<1xf32>, tensor<1xf32>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1xf32>, tensor<1xf32>)