func.func @main(%arg0: tensor<2xf32>) -> tensor<2xf32> {
  %0 = "tfl.custom"(%arg0) {custom_code = "stablehlo.rsqrt", custom_option = #tfl<const_bytes : "0x00000100002401">} : (tensor<2xf32>) -> tensor<2xf32>
  func.return %0 : tensor<2xf32>
}
}

// CHECK:       module
// CHECK-NEXT:  func @main(%arg0: tensor<2xf32>) -> tensor<2xf32> {
// CHECK-NEXT:  %0 = stablehlo.rsqrt %arg0 : tensor<2xf32>
// CHECK-NEXT:  return %0 : tensor<2xf32>
// CHECK-NEXT:  }

// tfl.log / tfl.logistic / tfl.max_pool_2d / tfl.mirror_pad / tfl.maximum /
// tfl.custom / tfl.mean / tfl.minimum / tfl.pad / tfl.pow / tfl.prelu /
// tfl.relu / tfl.relu6 / tfl.rsqrt / tfl.sin / tfl.slice / tfl.softmax /
// tfl.space_to_depth / tfl.sqrt / tfl.square / tfl.squared_difference /
// tfl.strided_slice / tfl.tanh / tfl.transpose / tfl.transpose_conv
class GpuBasicSupportedOpNoCost : public TargetHardwareOperation {

 Const 5
 GreaterEqual 1
 MatMul 1
 Mul 5
 Select 2
 Sigmoid 3
 Snapshot 1
 Split 1
 Tanh 2
cluster 29 size 9
 Add 1
 Mul 2
 RealDiv 2
 Sqrt 2
 Sub 2
cluster 30 size 9
 Add 1
 Mul 2
 RealDiv 2
 Sqrt 2
 Sub 2
cluster 31 size 4
 Mul 3
 UnsortedSegmentSum 1
cluster 32 size 4
 Mul 3

// and "bias" are found in src/math/bits.go

// Sqrt returns the square root of x.
//
// Special cases are:
//
//	Sqrt(+Inf) = +Inf
//	Sqrt(±0) = ±0
//	Sqrt(x < 0) = NaN
//	Sqrt(NaN) = NaN
func Sqrt(x float64) float64 {
	return sqrt(x)
}

// Note: On systems where Sqrt is a single instruction, the compiler
// may turn a direct call into a direct use of that instruction instead.

	sink64[4] = math.RoundToEven(x)
}

func sqrt(x float64) float64 {
	// amd64:"SQRTSD"
	// 386/sse2:"SQRTSD" 386/softfloat:-"SQRTD"
	// arm64:"FSQRTD"
	// arm/7:"SQRTD"
	// mips/hardfloat:"SQRTD" mips/softfloat:-"SQRTD"
	// mips64/hardfloat:"SQRTD" mips64/softfloat:-"SQRTD"
	// wasm:"F64Sqrt"
	// ppc64x:"FSQRT"
	// riscv64: "FSQRTD"
	return math.Sqrt(x)
}

func sqrt32(x float32) float32 {
	// amd64:"SQRTSS"

	}

	absx := Abs(x)

	var f float64
	var iu uint64
	k := 1
	if absx < Sqrt2M1 { //  |x| < Sqrt(2)-1
		if absx < Small { // |x| < 2**-29
			if absx < Tiny { // |x| < 2**-54
				return x
			}
			return x - x*x*0.5
		}
		if x > Sqrt2HalfM1 { // Sqrt(2)/2-1 < x
			// (Sqrt(2)/2-1) < x < (Sqrt(2)-1)
			k = 0
			f = x
			iu = 1
		}
	}
	var c float64
	if k != 0 {
		var u float64

//
//
// asinh(x)
// Method :
//	Based on
//	        asinh(x) = sign(x) * log [ |x| + sqrt(x*x+1) ]
//	we have
//	asinh(x) := x  if  1+x*x=1,
//	         := sign(x)*(log(x)+ln2) for large |x|, else
//	         := sign(x)*log(2|x|+1/(|x|+sqrt(x*x+1))) if|x|>2, else
//	         := sign(x)*log1p(|x| + x**2/(1 + sqrt(1+x**2)))
//

// Asinh returns the inverse hyperbolic sine of x.
//
// Special cases are:
//

// Use of this source code is governed by a BSD-style
// license that can be found in the LICENSE file.

package math

/*
	Hypot -- sqrt(p*p + q*q), but overflows only if the result does.
*/

// Hypot returns [Sqrt](p*p + q*q), taking care to avoid
// unnecessary overflow and underflow.
//
// Special cases are:
//
//	Hypot(±Inf, q) = +Inf
//	Hypot(p, ±Inf) = +Inf
//	Hypot(NaN, q) = NaN

	// standard deviation is sqrt( average( (x - xbar)^2 ) )
	// = sqrt( Integral( x^2 + xbar^2 -2*x*xbar dt ) / Duration )
	// = sqrt( ( Integral( x^2 dt ) + Duration * xbar^2 - 2*xbar*Integral(x dt) ) / Duration)
	// = sqrt( Integral(x^2 dt)/Duration - xbar^2 )
	variance := igr.IntegralXX/igr.ElapsedSeconds - avg*avg
	if variance >= 0 {
		return avg, math.Sqrt(variance)
	}
	return avg, math.NaN()

with an architecture tag, followed by a colon and a quoted Go-style
regexp to be matched. For example, the following test:

  func Sqrt(x float64) float64 {
  	   // amd64:"SQRTSD"
  	   // arm64:"FSQRTD"
  	   return math.Sqrt(x)
  }

verifies that math.Sqrt calls are intrinsified to a SQRTSD instruction
on amd64, and to a FSQRTD instruction on arm64.