x32 := math.Float32bits(f32)
	y32 := math.Float32bits(g32)
	if x32 != y32 {
		t.Errorf("got %x, want %x (diff=%x)", x32, y32, x32^y32)
	}

	f64 := math.Float64frombits(snan64bits)
	g64 := math.Float64frombits(snan64bitsVar)
	x64 := math.Float64bits(f64)
	y64 := math.Float64bits(g64)
	if x64 != y64 {
		t.Errorf("got %x, want %x (diff=%x)", x64, y64, x64^y64)
	}
}

func.func @Reciprocal_complexf64(%arg0: tensor<*xcomplex<f64>>) -> tensor<*xcomplex<f64>> {
  // CHECK: %[[ONE:.*]] = "tf.Const"() <{value = dense<(1.000000e+00,0.000000e+00)> : tensor<complex<f64>>}> : () -> tensor<complex<f64>>
  // CHECK: "tf.Div"(%[[ONE]], %arg0) : (tensor<complex<f64>>, tensor<*xcomplex<f64>>) -> tensor<*xcomplex<f64>>
  %0 = "tf.Reciprocal"(%arg0) : (tensor<*xcomplex<f64>>) -> tensor<*xcomplex<f64>>
  func.return %0 : tensor<*xcomplex<f64>>
}

			p.From.Offset = 0
		}

	case AMOVD:
		if p.From.Type == obj.TYPE_FCONST {
			f64 := p.From.Val.(float64)
			if math.Float64bits(f64) == 0 {
				p.As = AMOVV
				p.From.Type = obj.TYPE_REG
				p.From.Reg = REGZERO
				break
			}
			p.From.Type = obj.TYPE_MEM
			p.From.Sym = ctxt.Float64Sym(f64)
			p.From.Name = obj.NAME_EXTERN
			p.From.Offset = 0
		}
	}

	// Rewrite SUB constants into ADD.

// CHECK-LABEL: lower_rfft_to_rfft2d
func.func @lower_rfft_to_rfft2d(%input: tensor<10x20x30xf32>, %fft_len: tensor<1xi32>) -> tensor<10x20x30xcomplex<f64>> {
  %0 = "tf.RFFT"(%input, %fft_len) : (tensor<10x20x30xf32>, tensor<1xi32>) -> tensor<10x20x30xcomplex<f64>>
  func.return %0: tensor<10x20x30xcomplex<f64>>

// CHECK-DAG:  %[[CST:.*]] = arith.constant dense<-2> : tensor<i32>
// CHECK-DAG:  %[[CST0:.*]] = arith.constant dense<1> : tensor<1xi32>

  let arguments = (ins
     TFL_TensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$input,
     TFL_TensorOf<[I32]>:$dilations,
     TFL_0DTensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$padding_value
  );

  let results = (outs TFL_TensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$output);
}

def TFL_AddOp : TFL_Op<"add", [

    ) -> (tensor<i64>, tensor<i32>) {
    // CHECK:      %[[ANON_ITER:[0-9]*]] = "tf.AnonymousIteratorV3"
    // CHECK-SAME: output_shapes = [#tf_type.shape<32>, #tf_type.shape<64>]
    // CHECK-SAME: output_types = [f32, f64]
    // CHECK-NEXT: "tf.MakeIterator"(%[[ARG_0]], %[[ANON_ITER]])
    // CHECK-NEXT: %[[COND:.*]] = "tf.Const"
    // CHECK-NEXT: %[[WHILE:[0-9]*]]:5 = "tf.WhileRegion"(%[[COND]], %[[ARG_1]], %[[ARG_2]], %[[ARG_3]], %[[ARG_4]])

			return fmt.Errorf("converting NULL to %s is unsupported", dv.Kind())
		}
		s := asString(src)
		f64, err := strconv.ParseFloat(s, dv.Type().Bits())
		if err != nil {
			err = strconvErr(err)
			return fmt.Errorf("converting driver.Value type %T (%q) to a %s: %v", src, s, dv.Kind(), err)
		}
		dv.SetFloat(f64)
		return nil
	case reflect.String:
		if src == nil {

			p.From.Name = obj.NAME_EXTERN
			p.From.Offset = 0
		}

	case AFMOVD:
		if p.From.Type == obj.TYPE_FCONST {
			f64 := p.From.Val.(float64)
			if math.Float64bits(f64) == 0 { // +0
				break
			}
			p.From.Type = obj.TYPE_MEM
			p.From.Sym = ctxt.Float64Sym(f64)
			p.From.Name = obj.NAME_EXTERN
			p.From.Offset = 0
		}

		// put constants not loadable by LOAD IMMEDIATE into memory
	case AMOVD:

        tensor<?xf32>) -> tensor<?x?x?x?xf32>
  func.return %0 : tensor<?x?x?x?xf32>
}

// CHECK-LABEL: @batchNormInference_f64
// Validate that epsilon is properly promoted to f64
// CHECK-DAG: %[[EPS:.+]] = mhlo.constant dense<1.000000e+00> : tensor<256xf64>
func.func @batchNormInference_f64(
    %x: tensor<4x256xf64>, %scale: tensor<256xf64>, %offset: tensor<256xf64>,

  %0 = "tf.AddV2"(%arg1, %arg0) : (tensor<complex<f64>>, tensor<complex<f64>>) -> tensor<complex<f64>>
  func.return %0 : tensor<complex<f64>>
}

// -----

// CHECK-LABEL: @xla_variadic_reduce_v2_dynamic