}{
		{4, 4, true, 0, 0xFFFFFFF0},
		{4, 8, true, 28, 0x0FFFFFF0},
		{0, 0, true, 0, 0xFFFFFFFF},
		{8, 4, false, 0, 0},
		{0, 32, false, 0, 0},
		{0, 31, true, 1, 0x1},
		{31, 31, true, 0, 0x80000000},
		{32, 32, false, 0, 0},
	}
	for i, v := range tests {
		result := mergePPC64SldiSrw(v.sld, v.srw)
		if v.valid && result == 0 {
			t.Errorf("mergePPC64SldiSrw(Test %d) did not merge", i)
		} else if !v.valid && result != 0 {

//
//	for {
//		overflow = load(&b.overflow)
//		if uint32(overflow) == 0 {
//			// no pending entry
//			break
//		}
//		time = load(&b.overflowTime)
//		if cas(&b.overflow, overflow, ((overflow>>32)+1)<<32) {
//			// pending entry cleared
//			break
//		}
//	}
//	if uint32(overflow) > 0 {
//		emit entry for uint32(overflow), time
//	}
type profBuf struct {
	// accessed atomically
	r, w         profAtomic

  %2 = mhlo.constant dense<[1, 1, 1, 2, 2, 3, 3, 4, 4, 5, 5, 6, 6, 7, 7, 8, 8, 9, 9, 10, 10, 11, 11, 12, 12, 13, 13, 14, 14, 15, 15, 15]> : tensor<32xi64>
  %3 = mhlo.constant dense<16> : tensor<i64>

		{name: "SRDconst", argLength: 1, reg: gp11, asm: "SRD", aux: "Int64"},       // unsigned arg0 >> auxInt, 0 <= auxInt < 64, 64 bit width
		{name: "SRWconst", argLength: 1, reg: gp11, asm: "SRW", aux: "Int64"},       // unsigned arg0 >> auxInt, 0 <= auxInt < 32, 32 bit width

var CaseRanges = _CaseRanges
var _CaseRanges = []CaseRange{
	{0x0041, 0x005A, d{0, 32, 0}},
	{0x0061, 0x007A, d{-32, 0, -32}},
	{0x00B5, 0x00B5, d{743, 0, 743}},
	{0x00C0, 0x00D6, d{0, 32, 0}},
	{0x00D8, 0x00DE, d{0, 32, 0}},
	{0x00E0, 0x00F6, d{-32, 0, -32}},
	{0x00F8, 0x00FE, d{-32, 0, -32}},
	{0x00FF, 0x00FF, d{121, 0, 121}},
	{0x0100, 0x012F, d{UpperLower, UpperLower, UpperLower}},

	CiphertextSize       = k*encodingSize10 + encodingSize4
	EncapsulationKeySize = encryptionKeySize
	DecapsulationKeySize = decryptionKeySize + encryptionKeySize + 32 + 32
	SharedKeySize        = 32
	SeedSize             = 32 + 32
)

// A DecapsulationKey is the secret key used to decapsulate a shared key from a
// ciphertext. It includes various precomputed values.
type DecapsulationKey struct {

	SSECustomerKeySize = 32

	// SSEIVSize is the size of the IV data
	SSEIVSize = 32 // 32 bytes

	// SSEDAREPackageBlockSize - SSE dare package block size.
	SSEDAREPackageBlockSize = 64 * 1024 // 64KiB bytes

	// SSEDAREPackageMetaSize - SSE dare package meta padding bytes.
	SSEDAREPackageMetaSize = 32 // 32 bytes

)

// KMSKeyID returns in AWS compatible KMS KeyID() format.

// CHECK-SAME:                       %[[VAL_1:.*]]: tensor<32xi32>) -> tensor<1x32x10x32xi32> {
// CHECK:           %[[VAL_2:.*]] = "tf.AddV2"(%[[VAL_0]], %[[VAL_1]]) : (tensor<1x32x10x32xi32>, tensor<32xi32>) -> tensor<1x32x10x32xi32>
// CHECK:           return %[[VAL_2]] : tensor<1x32x10x32xi32>
// CHECK:         }
func.func @biasAdd_NHWC(%arg0: tensor<1x32x10x32xi32>, %arg1: tensor<32xi32>) -> tensor<1x32x10x32xi32> {

func.func @biasAdd(%arg0: tensor<1x10x10x32xf32>, %arg1: tensor<32xf32>) -> tensor<1x10x10x32xf32> {
  %0 = "tf.BiasAdd"(%arg0, %arg1) {T = "tfdtype$DT_FLOAT", data_format = "NHWC"} : (tensor<1x10x10x32xf32>, tensor<32xf32>) -> tensor<1x10x10x32xf32>
  func.return %0 : tensor<1x10x10x32xf32>

// CHECK-LABEL: biasAdd
// CHECK: tfl.add(%arg0, %arg1) <{fused_activation_function = "NONE"}> : (tensor<1x10x10x32xf32>, tensor<32xf32>) -> tensor<1x10x10x32xf32>
}

DATA ·IEEEConst+3544(SB)/8,$0x0000000074fe2232

	/* x^33856 mod p(x), x^33792 mod p(x) */
DATA ·IEEEConst+3552(SB)/8,$0x00000001e599099a
DATA ·IEEEConst+3560(SB)/8,$0x000000001ac6e2ba

	/* x^32832 mod p(x), x^32768 mod p(x) */
DATA ·IEEEConst+3568(SB)/8,$0x00000000fe128194
DATA ·IEEEConst+3576(SB)/8,$0x0000000013fca91e

	/* x^31808 mod p(x), x^31744 mod p(x) */
DATA ·IEEEConst+3584(SB)/8,$0x0000000077e8b990