func.return %dq : tensor<3x2xf32>
}

// -----

// CHECK-LABEL: @remove_volatile_qdq_with_requantization
// CHECK-SAME: %[[ARG0:.*]]: tensor<3x2xf32>
func.func @remove_volatile_qdq_with_requantization(%arg0: tensor<3x2xf32>) -> tensor<3x2xf32> {
  // CHECK: %[[Q1:.*]] = stablehlo.uniform_quantize %[[ARG0]]
  // CHECK: %[[Q2:.*]] = stablehlo.uniform_quantize %[[Q1]]
  // CHECK: %[[ABS:.*]] = stablehlo.abs %[[Q2]]

// CHECK-LABEL: identity
// CHECK: %0 = "tf.Identity"(%arg0) : (tensor<10xi32>) -> tensor<*xi32>
// CHECK: return %arg0, %arg1, %arg2, %0
}

func.func @sharding(%arg0: tensor<10x10xi32>) -> (tensor<10x10xi32>) {
  %0 = "tf.MatMul"(%arg0, %arg0) {device = "", transpose_a = false, transpose_b = false} : (tensor<10x10xi32>, tensor<10x10xi32>) -> tensor<10x10xi32>

func.func @FuseTransposeFCRhsToBatchMatmul(%arg0: tensor<16x1024xf32>, %arg1: tensor<1024x128xf32>, %arg2: none) -> tensor<16x128xf32> {
  %cst = arith.constant dense<[1, 0]> : tensor<2xi32>
  %0 = "tfl.transpose"(%arg1, %cst) : (tensor<1024x128xf32>, tensor<2xi32>) -> tensor<128x1024xf32>
  // CHECK: "tfl.batch_matmul"(%arg0, %arg1)

func (m *MockRuntimeCache) ForceUpdateIfOlder(arg0 context.Context, arg1 time.Time) error {
	m.ctrl.T.Helper()
	ret := m.ctrl.Call(m, "ForceUpdateIfOlder", arg0, arg1)
	ret0, _ := ret[0].(error)
	return ret0
}

// ForceUpdateIfOlder indicates an expected call of ForceUpdateIfOlder.
func (mr *MockRuntimeCacheMockRecorder) ForceUpdateIfOlder(arg0, arg1 any) *gomock.Call {
	mr.mock.ctrl.T.Helper()

// CHECK: %[[dw:.*]] = "tfl.depthwise_conv_2d"(%arg0, %arg1, %arg2)
// CHECK: %[[q:.*]] = "tfl.quantize"(%[[dw]])
// CHECK: %[[dq:.*]] = "tfl.dequantize"(%[[q]])
// CHECK: return %[[dq]], %[[dq]]
}

// Series of values needing requantization -- first the args then the results
// of concatenation operations. concat(concat(arg2, arg0), concat(arg1, arg0)),
// concat(concat(arg2, arg0), arg3)). arg0 should be requantized twice --

//go:build msan

#include "go_asm.h"
#include "textflag.h"

#define RARG0 R0
#define RARG1 R1
#define RARG2 R2
#define FARG R3

// func runtime·domsanread(addr unsafe.Pointer, sz uintptr)
// Called from msanread.
TEXT	runtime·domsanread(SB), NOSPLIT, $0-16
	MOVD	addr+0(FP), RARG0
	MOVD	sz+8(FP), RARG1
	// void __msan_read_go(void *addr, uintptr_t sz);
	MOVD	$__msan_read_go(SB), FARG

  %0 = "tfl.maximum"(%arg0, %cst) : (tensor<2x3xbf16>, tensor<2x3xbf16>) -> tensor<2x3xbf16>
  func.return %0 : tensor<2x3xbf16>

  // CHECK: %[[RESULT:.*]] = "tfl.relu"(%arg0)
  // CHECK: return %[[RESULT]]
}

// CHECK-LABEL: fuseScalarAddIntoConv2dBf16
func.func @fuseScalarAddIntoConv2dBf16(%arg0: tensor<256x32x32x3xbf16>, %arg1: tensor<16x3x3x3xbf16>) -> tensor<256x8x7x16xbf16> {
  %cst = arith.constant dense<1.5> : tensor<bf16>

// CHECK-LABEL: func @testPwStreamResults
func.func @testPwStreamResults(%arg0: tensor<f32>, %arg1: tensor<f32>) {
  "tf.PwStreamResults"(%arg0, %arg1) {names = ["foo", "bar"]} : (tensor<f32>, tensor<f32>) -> ()
  return
}

// -----
// CHECK-LABEL: func @test_ifrt_call
func.func @test_ifrt_call(%arg0: tensor<?xf32>, %arg1: tensor<?xf32>) {

  // CHECK-LABEL: @serving_default
  // CHECK-SAME: %[[ARG0:[^:[:space:]]+]]
  // CHECK-SAME: %[[ARG1:[^:[:space:]]+]]
  func.func private @serving_default(%arg0: tensor<1x4xf32>, %arg1: tensor<4x3xf32>) -> tensor<1x3xf32> {

// CHECK-NEXT: z.b
// CHECK-NEXT: z.c

// -----

// CHECK-LABEL: @terminator
func.func @terminator(%arg0: f32) -> (f32) attributes {ignore_side_effects_for_testing} {
  func.call @terminator(%arg0) : (f32) -> (f32)
  "x.a"(%arg0) : (f32) -> ()
  "y.a"(%arg0) : (f32) -> ()
  "z.a"(%arg0) : (f32) -> ()
  func.return %arg0 : f32
}
// CHECK: x.a
// CHECK: y.a
// CHECK: z.a
// CHECK: return
// CHECK-NOT: call
// CHECK: }