//       VADDR=0xffffffff81000000
		// stextOffset=0xffffffff81000198
		return start - *stextOffset, true
	}
	return 0, false
}

// GetBase determines the base address to subtract from virtual
// address to get symbol table address. For an executable, the base
// is 0. Otherwise, it's a shared library, and the base is the
// address where the mapping starts. The kernel needs special handling.

      %cst = stablehlo.constant dense<1.0> : tensor<1x3xf32>
      %0 = stablehlo.dot_general %arg0, %arg1, contracting_dims = [1] x [0], precision = [] : (tensor<1x1024xf32>, tensor<1024x3xf32>) -> tensor<1x3xf32>
      %1 = stablehlo.subtract %cst, %0 : tensor<1x3xf32>
      %2 = stablehlo.add %0, %cst : tensor<1x3xf32>
      return %2 : tensor<1x3xf32>
    }
  }
)mlir";

constexpr absl::string_view kModuleXlaCallModule = R"mlir(
  module {

    auto matmul_result =
        BatchMatMul(scope, grad_grad_expand, logits_softmax_expand);
    axis.push_back(1);
    auto squeeze_result = Squeeze(scope, matmul_result, Squeeze::Axis(axis));
    auto subtraction_result = Subtract(scope, grad_grad, squeeze_result);
    auto multiply_result = Multiply(scope, subtraction_result, logits_softmax);
    grad = Add(scope, grad, multiply_result);
  }

    @Override
    BigInteger toX(double d, RoundingMode mode) {
      return DoubleMath.roundToBigInteger(d, mode);
    }

    @Override
    BigInteger minus(BigInteger a, BigInteger b) {
      return a.subtract(b);
    }
  }

  /**
   * Returns the result of dividing {@code p} by {@code q}, rounding using the specified {@code
   * RoundingMode}.
   *

    @Override
    BigInteger toX(double d, RoundingMode mode) {
      return DoubleMath.roundToBigInteger(d, mode);
    }

    @Override
    BigInteger minus(BigInteger a, BigInteger b) {
      return a.subtract(b);
    }
  }

  /**
   * Returns the result of dividing {@code p} by {@code q}, rounding using the specified {@code
   * RoundingMode}.
   *

  func.return %0: tensor<2xi32>
}

// CHECK-LABEL: func @sub
func.func @sub(%arg0: tensor<2xi32>) -> tensor<2xi32> {
  // CHECK-NEXT:  %0 = mhlo.subtract %arg0, %arg0 : tensor<2xi32>
  // CHECK-NEXT:  return %0 : tensor<2xi32>
  %0 = "tf.Sub"(%arg0, %arg0) : (tensor<2xi32>, tensor<2xi32>) -> tensor<2xi32>
  func.return %0: tensor<2xi32>
}

			mheap_.central[spc].mcentral.fullSwept(sweepgen).push(s)
			return false
		}

		// It's only at this point that the sweeper doesn't actually need to look
		// at this arena anymore, so subtract from pagesInUse now.
		mheap_.pagesInUse.Add(-s.npages)
		s.state.set(mSpanDead)

		// The arena is ready to be recycled. Remove it from the quarantine list

  %0 = "mhlo.broadcast_in_dim"(%arg0) <{broadcast_dimensions = dense<[0, 1]> : tensor<2xi64>}> : (tensor<1x1xf32>) -> tensor<1x1000xf32>
  %1 = mhlo.subtract %0, %arg1 : tensor<1x1000xf32>
  %2 = mhlo.subtract %arg1, %0 : tensor<1x1000xf32>
  func.return %1, %2 : tensor<1x1000xf32>, tensor<1x1000xf32>
}

// CHECK-LABEL:   func @broadcast_sub_chlo(

            dri.fixupHost(server);
        }

        if ( log.isDebugEnabled() ) {
            log.debug("Adding key " + key + " to " + dr);
        }

        /*
         * Subtract the server and share from the pathConsumed so that
         * it reflects the part of the relative path consumed and not
         * the entire path.
         */

		sa := new(SockaddrUnix)

		// Some versions of AIX have a bug in getsockname (see IV78655).
		// We can't rely on sa.Len being set correctly.
		n := SizeofSockaddrUnix - 3 // subtract leading Family, Len, terminating NUL.
		for i := 0; i < n; i++ {
			if pp.Path[i] == 0 {
				n = i
				break
			}
		}
		sa.Name = string(unsafe.Slice((*byte)(unsafe.Pointer(&pp.Path[0])), n))