j += 2
		fmt.Fprintf(w, "\n")
	}

	for i := (1024 * 2) - 128; i >= 0; i -= 128 {
		a := reflect_bits(get_remainder(ref_poly, 32, uint(i+32)), 32)
		b := reflect_bits(get_remainder(ref_poly, 32, uint(i+64)), 32)
		c := reflect_bits(get_remainder(ref_poly, 32, uint(i+96)), 32)
		d := reflect_bits(get_remainder(ref_poly, 32, uint(i+128)), 32)

  %0 = mhlo.constant dense<0xFF800000> : tensor<f32>
  %1 = mhlo.constant dense<0> : tensor<i32>
  %2 = "mhlo.iota"() <{iota_dimension = 0 : i64}> : () -> tensor<32xi32>
  %3 = "mhlo.reshape"(%2) : (tensor<32xi32>) -> tensor<1x32x1xi32>
  %4:2 = mhlo.reduce(%arg0 init: %0), (%3 init: %1) across dimensions = [1] : (tensor<1x32x1xf32>, tensor<1x32x1xi32>, tensor<f32>, tensor<i32>) -> (tensor<1x1xf32>, tensor<1x1xi32>)

  %2 = "tfl.pseudo_qconst"() {qtype = tensor<32x!quant.uniform<i32:f32, 1.7052092479439231E-4>>, value = dense<0> : tensor<32xi32>} : () -> tensor<32x!quant.uniform<i32:f32, 1.7052092479439231E-4>>

//
//	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> {