// CHECK: %[[ADD1:.*]] = "tf.AddV2"(%[[OLD_SLICE1]], %[[UPDATE_SLICE1]])
  // CHECK: %[[UPDATE1:.*]] = "tf.XlaDynamicUpdateSlice"(%[[UPDATE0]], %[[ADD1]]
  // CHECK-SAME: (tensor<5x3xf32>, tensor<1x3xf32>, tensor<2xi32>) -> tensor<5x3xf32>
  // CHECK: "tf.AssignVariableOp"(%[[VAR]], %[[UPDATE1]])

  %4 = "tf.Log"(%3) {device = "/job:localhost/replica:0/task:0/device:GPU:0"}: (tensor<4x4xf32>) -> tensor<4x4xf32>

  // CHECK: %[[ADD1:.*]] = "tf.AddV2"
  // CHECK: %[[LOG1:.*]] = "tf.Log"(%[[ADD1]])
  %5 = "tf.AddV2"(%4, %1): (tensor<4x4xf32>, tensor<1xf32>) -> tensor<4x4xf32>
  %6 = "tf.Log"(%5): (tensor<4x4xf32>) -> tensor<4x4xf32>

        when:
        from.clear()

        then:
        from.empty
        collection.from.empty

        when:
        def add1 = from.add('a')
        def add2 = from.add('b')
        def add3 = from.add('a')

        then:
        add1
        add2
        !add3

        and:
        from as List == ['a', 'b']
        collection.from as List == ['a', 'b']

        when:

// CHECK-DAG:  %[[CONSTANT1:.*]] = arith.constant dense<0.000000e+00> : tensor<4xf32>
//              variance + epsilon
// CHECK:  %[[ADD1:.*]] = "tf.Add"(%[[ARG4:.*]], %[[CONSTANT]])
//              rsqrt(variance + epsilon)
// CHECK:  %[[RSQRT:.*]] = "tf.Rsqrt"(%[[ADD1]])
//              scale * rsqrt(variance + epsilon)
// CHECK:  %[[MUL1:.*]] = "tf.Mul"(%[[ARG1:.*]], %[[RSQRT]])

            %island1:3 = tf_executor.island {
              %add1 = "tf.Add"(%arg0, %arg1) : (tensor<*xi32>, tensor<i32>) -> tensor<*xi32>
              %add2 = "tf.Add"(%add1, %arg1) : (tensor<*xi32>, tensor<i32>) -> tensor<*xi32>
              %res = "tf.Print"(%add2) { message = "add result" } : (tensor<*xi32>) -> (tensor<*xi32>)
              tf_executor.yield %add1, %add2 : tensor<*xi32>, tensor<*xi32>
            }

        %island1:3 = tf_executor.island {
          %add1 = "tf.Add"(%arg0, %arg1) : (tensor<*xi32>, tensor<i32>) -> tensor<*xi32>
          %add2 = "tf.Add"(%add1, %arg1) : (tensor<*xi32>, tensor<i32>) -> tensor<*xi32>
          %res = "tf.Print"(%add2) { message = "add result" } : (tensor<*xi32>) -> (tensor<*xi32>)
          tf_executor.yield %add1, %add2 : tensor<*xi32>, tensor<*xi32>
        }

// Fusing:  %[[add:[0-9].*]] = tfl.add %arg0, %arg1 {fused_activation_function = "NONE"} : tensor<1xf32>
// Fusing:  %[[add1:[0-9].*]] = tfl.add %arg0, %[[add]] {fused_activation_function = "RELU"} : tensor<1xf32>
// Fusing:  %[[relu:[0-9].*]] = "tfl.relu"(%arg0) : (tensor<1xf32>) -> tensor<1xf32>
// Fusing:  %[[add2:[0-9].*]] = tfl.add %[[relu]], %[[add1]] {fused_activation_function = "RELU6"} : tensor<1xf32>

  }, {
    ^bb0(%barg0: tensor<*xf32>, %barg1: tensor<i32>):
      %add0 = "tf.Add"(%barg0, %barg0) : (tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32>
      %add1 = "tf.Add"(%barg1, %barg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
      "tf.Yield"(%add0, %add1) : (tensor<*xf32>, tensor<i32>) -> ()
  }) { is_stateless = false } : (tensor<*xf32>, tensor<i32>) -> (tensor<*xf32>, tensor<i32>)

  func.return %0#0 : tensor<*xf32>
}

openAVX2Tail128LoopB:
	ADDQ     $16, itr2
	chachaQR_AVX2(AA1, BB1, CC1, DD1, TT0)
	VPALIGNR $4, BB1, BB1, BB1
	VPALIGNR $8, CC1, CC1, CC1
	VPALIGNR $12, DD1, DD1, DD1
	chachaQR_AVX2(AA1, BB1, CC1, DD1, TT0)
	VPALIGNR $12, BB1, BB1, BB1
	VPALIGNR $8, CC1, CC1, CC1
	VPALIGNR $4, DD1, DD1, DD1
	CMPQ     itr2, itr1
	JB       openAVX2Tail128LoopA
	CMPQ     itr2, $160
	JNE      openAVX2Tail128LoopB

  // CHECK-DAG: %[[ADD1:.*]] = chlo.broadcast_add %[[FEATURES]], %[[ONE]] {broadcast_dimensions = array<i64>}
  // CHECK-DAG: %[[MULGRAD:.*]] = mhlo.multiply %[[GRADIENTS]], %[[ADD1]] : (tensor<4x8xf32>, tensor<?x?xf32>) -> tensor<4x8xf32>
  // CHECK: %[[RESULT:.*]] = mhlo.select %[[PRED]], %[[GRADIENTS]], %[[MULGRAD]]