OpenCL variable-amount rotate(x, y) extra mask instruction in PTX

Accelerated Computing GPU-Accelerated Libraries
1

NVidia Developers,

I have an OpenCL macro:

'define ROTL(x, s) rotate((uint)(x), (uint)(s))

It appears that the NVVM Compiler generates an explicit bitwise AND (and.b32) before every variable rotate / funnel shift (shf.l.wrap.b32) in the PTX compilation.

This type of sequence repeats over and over when the funnel shift is variable:

and.b32 %r37, %r32, 31;
shf.l.wrap.b32 %r38, %r36, %r36, %r37;

According to the PTX ISA documention semantics, the funnel shift already masks the shift amount to 5 bits, so my thinking is this extra mask instruction is unnecessary.

I have experimented with inlining the funnel shift asm directly in lieu of using the OpenCL built in rotate(x, y) and the compiler does omit the and.b32 and the code works as expected. In the case of this algorithm, which is close to 700 instructions, almost 100 of them are unnecessary if rotate (x, y) is used.

I am not sure if this is a bug report, an optimization request or if I am misunderstanding something about OpenCL rotate(x, y) or NVidia’s funnel shift. I am looking for some feedback.

I am using the Linux 595.71.05 driver and targeted sm_75.

OpenCL kernel, PTX and SASS are attached as a zip.

rc5-72.zip (14.4 KB)

Thank you and please advise!

4 
//CORENAME=ocl_rc572_1pipe_src
#define ROTL(x, s) rotate((uint)(x), (uint)(s))
#define ROTL3(x)   rotate((uint)(x), 3u)
#define SWAP(x)    (as_uint(as_uchar4((x)).wzyx))

#define P 0xB7E15163
#define Q 0x9E3779B9

#define ROUND1(a, b, c, d) \
  S[a] = ROTL3(S[b] + P + a*Q + L[c]); \
  t = S[a] + L[c]; \
  L[d] = ROTL(L[d] + t, t)

#define ROUND23(a, b, c, d) \
  S[a] = ROTL3(S[a] + S[b] + L[c]); \
  t = S[a] + L[c]; \
  L[d] = ROTL(L[d] + t, t)

#define ENCRYPT(a) \
    A = ROTL(A^B, B) + S[a]; \
    B = ROTL(B^A, A) + S[a+1]

__kernel void ocl_rc572_1pipe( __constant uint *rc5_72unitwork, __global uint *outbuf)
{
  uint L[3];
  uint S[26];
  uint A,B;
  uint t;

  L[2] = rc5_72unitwork[0];   //L0hi;
  L[1] = rc5_72unitwork[1];   //L0mid;
  L[0] = rc5_72unitwork[8];   
  S[1] = rc5_72unitwork[9];

  L[2] += get_global_id(0);
  uint l1_t1 = L[1];
  uint l1_t2 = l1_t1 + (L[2] >> 8);
  L[2] &= 0x000000ff;
  if(l1_t2 < l1_t1)
  {
    uint l0_t = SWAP(rc5_72unitwork[2]);
    l0_t +=1;
    L[0] = ROTL(0xBF0A8B1D + SWAP(l0_t), 0x1d);
    S[1] = ROTL3(L[0] + 0xBF0A8B1D + 0x5618cb1c);
  }
  L[1] = SWAP(l1_t2);
  
  S[0] = 0xBF0A8B1D;
  t = S[1] + L[0];
  L[1] = ROTL(L[1] + t, t);

  ROUND1( 2,  1, 1, 2);
  ROUND1( 3,  2, 2, 0);
  ROUND1( 4,  3, 0, 1);
  ROUND1( 5,  4, 1, 2);
  ROUND1( 6,  5, 2, 0);
  ROUND1( 7,  6, 0, 1);
  ROUND1( 8,  7, 1, 2);
  ROUND1( 9,  8, 2, 0);
  ROUND1(10,  9, 0, 1);
  ROUND1(11, 10, 1, 2);
  ROUND1(12, 11, 2, 0);
  ROUND1(13, 12, 0, 1);
  ROUND1(14, 13, 1, 2);
  ROUND1(15, 14, 2, 0);
  ROUND1(16, 15, 0, 1);
  ROUND1(17, 16, 1, 2);
  ROUND1(18, 17, 2, 0);
  ROUND1(19, 18, 0, 1);
  ROUND1(20, 19, 1, 2);
  ROUND1(21, 20, 2, 0);
  ROUND1(22, 21, 0, 1);
  ROUND1(23, 22, 1, 2);
  ROUND1(24, 23, 2, 0);
  ROUND1(25, 24, 0, 1);

  ROUND23(0, 25, 1, 2);
  ROUND23(1,  0, 2, 0);
  ROUND23(2,  1, 0, 1);
  ROUND23(3,  2, 1, 2);
  ROUND23(4,  3, 2, 0);
  ROUND23(5,  4, 0, 1);
  ROUND23(6,  5, 1, 2);
  ROUND23(7,  6, 2, 0);
  ROUND23(8,  7, 0, 1);
  ROUND23(9,  8, 1, 2);
  ROUND23(10, 9, 2, 0);
  ROUND23(11,10, 0, 1);
  ROUND23(12,11, 1, 2);
  ROUND23(13,12, 2, 0);
  ROUND23(14,13, 0, 1);
  ROUND23(15,14, 1, 2);
  ROUND23(16,15, 2, 0);
  ROUND23(17,16, 0, 1);
  ROUND23(18,17, 1, 2);
  ROUND23(19,18, 2, 0);
  ROUND23(20,19, 0, 1);
  ROUND23(21,20, 1, 2);
  ROUND23(22,21, 2, 0);
  ROUND23(23,22, 0, 1);
  ROUND23(24,23, 1, 2);
  ROUND23(25,24, 2, 0);

  ROUND23(0, 25, 0, 1);
  ROUND23(1,  0, 1, 2);
  ROUND23(2,  1, 2, 0);
  ROUND23(3,  2, 0, 1);
  ROUND23(4,  3, 1, 2);
  ROUND23(5,  4, 2, 0);
  ROUND23(6,  5, 0, 1);
  ROUND23(7,  6, 1, 2);
  ROUND23(8,  7, 2, 0);
  ROUND23(9,  8, 0, 1);
  ROUND23(10, 9, 1, 2);
  ROUND23(11,10, 2, 0);
  ROUND23(12,11, 0, 1);
  ROUND23(13,12, 1, 2);
  ROUND23(14,13, 2, 0);
  ROUND23(15,14, 0, 1);
  ROUND23(16,15, 1, 2);
  ROUND23(17,16, 2, 0);
  ROUND23(18,17, 0, 1);
  ROUND23(19,18, 1, 2);
  ROUND23(20,19, 2, 0);
  ROUND23(21,20, 0, 1);
  ROUND23(22,21, 1, 2);
  ROUND23(23,22, 2, 0);

  S[24] = ROTL3(S[24] + S[23] + L[0]); 

  A = rc5_72unitwork[4] + S[0];	//plain_lo
  B = rc5_72unitwork[5] + S[1]; //plain_hi

  ENCRYPT(2);
  ENCRYPT(4);
  ENCRYPT(6);
  ENCRYPT(8);
  ENCRYPT(10);
  ENCRYPT(12);
  ENCRYPT(14);
  ENCRYPT(16);
  ENCRYPT(18);
  ENCRYPT(20);
  ENCRYPT(22);

  A = ROTL(A^B,B) + S[24]; 

  if(A == rc5_72unitwork[6])
  {
    t = S[24] + L[0]; 
    L[1] = ROTL(L[1] + t, t);

    S[25] = ROTL3(S[25] + S[24] +L[1]);
    B = ROTL(B^A,A)+S[25];

    uint idx = atomic_add(&outbuf[0], 1)*2+1;
    uint val = get_global_id(0) + rc5_72unitwork[3]; //keyN+offset
    uint attrib = (B == rc5_72unitwork[7])?0x80000000:0;
    outbuf[idx] = attrib;
    outbuf[idx+1] = val;
  }
}
5 
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-37855471
// Unknown Toolkit Version
// Based on NVVM 7.0.1
//

.version 9.2
.target sm_75, texmode_independent
.address_size 64

	// .globl	ocl_rc572_1pipe

.entry ocl_rc572_1pipe(
	.param .u64 .ptr .const .align 4 ocl_rc572_1pipe_param_0,
	.param .u64 .ptr .global .align 4 ocl_rc572_1pipe_param_1
)
{
	.reg .pred 	%p<4>;
	.reg .b32 	%r<679>;
	.reg .b64 	%rd<10>;


	ld.param.u64 	%rd2, [ocl_rc572_1pipe_param_0];
	ld.const.u32 	%r16, [%rd2];
	ld.const.u32 	%r677, [%rd2+32];
	ld.const.u32 	%r678, [%rd2+36];
	mov.u32 	%r17, %ctaid.x;
	mov.u32 	%r18, %ntid.x;
	mov.u32 	%r19, %tid.x;
	mov.b32 	%r20, %envreg3;
	add.s32 	%r21, %r19, %r20;
	mad.lo.s32 	%r3, %r18, %r17, %r21;
	add.s32 	%r4, %r16, %r3;
	shr.u32 	%r22, %r4, 8;
	ld.const.u32 	%r23, [%rd2+4];
	add.s32 	%r5, %r22, %r23;
	setp.ge.u32 	%p1, %r5, %r22;
	@%p1 bra 	$L__BB0_2;

	ld.const.u32 	%r24, [%rd2+8];
	mov.u32 	%r25, 291;
	mov.u32 	%r26, 0;
	prmt.b32 	%r27, %r24, %r26, %r25;
	add.s32 	%r28, %r27, 1;
	prmt.b32 	%r29, %r28, %r26, %r25;
	add.s32 	%r30, %r29, -1089828067;
	shf.l.wrap.b32 	%r677, %r30, %r30, 29;
	add.s32 	%r31, %r677, 354637369;
	shf.l.wrap.b32 	%r678, %r31, %r31, 3;

$L__BB0_2:
	add.s32 	%r32, %r678, %r677;
	mov.u32 	%r33, 291;
	mov.u32 	%r34, 0;
	prmt.b32 	%r35, %r5, %r34, %r33;
	add.s32 	%r36, %r32, %r35;
	and.b32  	%r37, %r32, 31;
	shf.l.wrap.b32 	%r38, %r36, %r36, %r37;
	add.s32 	%r39, %r678, %r38;
	add.s32 	%r40, %r39, -196066091;
	shf.l.wrap.b32 	%r41, %r40, %r40, 3;
	add.s32 	%r42, %r41, %r38;
	and.b32  	%r43, %r4, 255;
	add.s32 	%r44, %r42, %r43;
	and.b32  	%r45, %r42, 31;
	shf.l.wrap.b32 	%r46, %r44, %r44, %r45;
	add.s32 	%r47, %r41, %r46;
	add.s32 	%r48, %r47, -1836597618;
	shf.l.wrap.b32 	%r49, %r48, %r48, 3;
	add.s32 	%r50, %r49, %r46;
	add.s32 	%r51, %r50, %r677;
	and.b32  	%r52, %r50, 31;
	shf.l.wrap.b32 	%r53, %r51, %r51, %r52;
	add.s32 	%r54, %r49, %r53;
	add.s32 	%r55, %r54, 817838151;
	shf.l.wrap.b32 	%r56, %r55, %r55, 3;
	add.s32 	%r57, %r56, %r53;
	add.s32 	%r58, %r57, %r38;
	and.b32  	%r59, %r57, 31;
	shf.l.wrap.b32 	%r60, %r58, %r58, %r59;
	add.s32 	%r61, %r56, %r60;
	add.s32 	%r62, %r61, -822693376;
	shf.l.wrap.b32 	%r63, %r62, %r62, 3;
	add.s32 	%r64, %r63, %r60;
	add.s32 	%r65, %r64, %r46;
	and.b32  	%r66, %r64, 31;
	shf.l.wrap.b32 	%r67, %r65, %r65, %r66;
	add.s32 	%r68, %r63, %r67;
	add.s32 	%r69, %r68, 1831742393;
	shf.l.wrap.b32 	%r70, %r69, %r69, 3;
	add.s32 	%r71, %r70, %r67;
	add.s32 	%r72, %r71, %r53;
	and.b32  	%r73, %r71, 31;
	shf.l.wrap.b32 	%r74, %r72, %r72, %r73;
	add.s32 	%r75, %r70, %r74;
	add.s32 	%r76, %r75, 191210866;
	shf.l.wrap.b32 	%r77, %r76, %r76, 3;
	add.s32 	%r78, %r77, %r74;
	add.s32 	%r79, %r78, %r60;
	and.b32  	%r80, %r78, 31;
	shf.l.wrap.b32 	%r81, %r79, %r79, %r80;
	add.s32 	%r82, %r77, %r81;
	add.s32 	%r83, %r82, -1449320661;
	shf.l.wrap.b32 	%r84, %r83, %r83, 3;
	add.s32 	%r85, %r84, %r81;
	add.s32 	%r86, %r85, %r67;
	and.b32  	%r87, %r85, 31;
	shf.l.wrap.b32 	%r88, %r86, %r86, %r87;
	add.s32 	%r89, %r84, %r88;
	add.s32 	%r90, %r89, 1205115108;
	shf.l.wrap.b32 	%r91, %r90, %r90, 3;
	add.s32 	%r92, %r91, %r88;
	add.s32 	%r93, %r92, %r74;
	and.b32  	%r94, %r92, 31;
	shf.l.wrap.b32 	%r95, %r93, %r93, %r94;
	add.s32 	%r96, %r91, %r95;
	add.s32 	%r97, %r96, -435416419;
	shf.l.wrap.b32 	%r98, %r97, %r97, 3;
	add.s32 	%r99, %r98, %r95;
	add.s32 	%r100, %r99, %r81;
	and.b32  	%r101, %r99, 31;
	shf.l.wrap.b32 	%r102, %r100, %r100, %r101;
	add.s32 	%r103, %r98, %r102;
	add.s32 	%r104, %r103, -2075947946;
	shf.l.wrap.b32 	%r105, %r104, %r104, 3;
	add.s32 	%r106, %r105, %r102;
	add.s32 	%r107, %r106, %r88;
	and.b32  	%r108, %r106, 31;
	shf.l.wrap.b32 	%r109, %r107, %r107, %r108;
	add.s32 	%r110, %r105, %r109;
	add.s32 	%r111, %r110, 578487823;
	shf.l.wrap.b32 	%r112, %r111, %r111, 3;
	add.s32 	%r113, %r112, %r109;
	add.s32 	%r114, %r113, %r95;
	and.b32  	%r115, %r113, 31;
	shf.l.wrap.b32 	%r116, %r114, %r114, %r115;
	add.s32 	%r117, %r112, %r116;
	add.s32 	%r118, %r117, -1062043704;
	shf.l.wrap.b32 	%r119, %r118, %r118, 3;
	add.s32 	%r120, %r119, %r116;
	add.s32 	%r121, %r120, %r102;
	and.b32  	%r122, %r120, 31;
	shf.l.wrap.b32 	%r123, %r121, %r121, %r122;
	add.s32 	%r124, %r119, %r123;
	add.s32 	%r125, %r124, 1592392065;
	shf.l.wrap.b32 	%r126, %r125, %r125, 3;
	add.s32 	%r127, %r126, %r123;
	add.s32 	%r128, %r127, %r109;
	and.b32  	%r129, %r127, 31;
	shf.l.wrap.b32 	%r130, %r128, %r128, %r129;
	add.s32 	%r131, %r126, %r130;
	add.s32 	%r132, %r131, -48139462;
	shf.l.wrap.b32 	%r133, %r132, %r132, 3;
	add.s32 	%r134, %r133, %r130;
	add.s32 	%r135, %r134, %r116;
	and.b32  	%r136, %r134, 31;
	shf.l.wrap.b32 	%r137, %r135, %r135, %r136;
	add.s32 	%r138, %r133, %r137;
	add.s32 	%r139, %r138, -1688670989;
	shf.l.wrap.b32 	%r140, %r139, %r139, 3;
	add.s32 	%r141, %r140, %r137;
	add.s32 	%r142, %r141, %r123;
	and.b32  	%r143, %r141, 31;
	shf.l.wrap.b32 	%r144, %r142, %r142, %r143;
	add.s32 	%r145, %r140, %r144;
	add.s32 	%r146, %r145, 965764780;
	shf.l.wrap.b32 	%r147, %r146, %r146, 3;
	add.s32 	%r148, %r147, %r144;
	add.s32 	%r149, %r148, %r130;
	and.b32  	%r150, %r148, 31;
	shf.l.wrap.b32 	%r151, %r149, %r149, %r150;
	add.s32 	%r152, %r147, %r151;
	add.s32 	%r153, %r152, -674766747;
	shf.l.wrap.b32 	%r154, %r153, %r153, 3;
	add.s32 	%r155, %r154, %r151;
	add.s32 	%r156, %r155, %r137;
	and.b32  	%r157, %r155, 31;
	shf.l.wrap.b32 	%r158, %r156, %r156, %r157;
	add.s32 	%r159, %r154, %r158;
	add.s32 	%r160, %r159, 1979669022;
	shf.l.wrap.b32 	%r161, %r160, %r160, 3;
	add.s32 	%r162, %r161, %r158;
	add.s32 	%r163, %r162, %r144;
	and.b32  	%r164, %r162, 31;
	shf.l.wrap.b32 	%r165, %r163, %r163, %r164;
	add.s32 	%r166, %r161, %r165;
	add.s32 	%r167, %r166, 339137495;
	shf.l.wrap.b32 	%r168, %r167, %r167, 3;
	add.s32 	%r169, %r168, %r165;
	add.s32 	%r170, %r169, %r151;
	and.b32  	%r171, %r169, 31;
	shf.l.wrap.b32 	%r172, %r170, %r170, %r171;
	add.s32 	%r173, %r168, %r172;
	add.s32 	%r174, %r173, -1301394032;
	shf.l.wrap.b32 	%r175, %r174, %r174, 3;
	add.s32 	%r176, %r175, %r172;
	add.s32 	%r177, %r176, %r158;
	and.b32  	%r178, %r176, 31;
	shf.l.wrap.b32 	%r179, %r177, %r177, %r178;
	add.s32 	%r180, %r175, %r179;
	add.s32 	%r181, %r180, 1353041737;
	shf.l.wrap.b32 	%r182, %r181, %r181, 3;
	add.s32 	%r183, %r182, %r179;
	add.s32 	%r184, %r183, %r165;
	and.b32  	%r185, %r183, 31;
	shf.l.wrap.b32 	%r186, %r184, %r184, %r185;
	add.s32 	%r187, %r182, %r186;
	add.s32 	%r188, %r187, -287489790;
	shf.l.wrap.b32 	%r189, %r188, %r188, 3;
	add.s32 	%r190, %r189, %r186;
	add.s32 	%r191, %r190, %r172;
	and.b32  	%r192, %r190, 31;
	shf.l.wrap.b32 	%r193, %r191, %r191, %r192;
	add.s32 	%r194, %r189, %r193;
	add.s32 	%r195, %r194, -1928021317;
	shf.l.wrap.b32 	%r196, %r195, %r195, 3;
	add.s32 	%r197, %r196, %r193;
	add.s32 	%r198, %r197, %r179;
	and.b32  	%r199, %r197, 31;
	shf.l.wrap.b32 	%r200, %r198, %r198, %r199;
	add.s32 	%r201, %r196, %r200;
	add.s32 	%r202, %r201, 726414452;
	shf.l.wrap.b32 	%r203, %r202, %r202, 3;
	add.s32 	%r204, %r203, %r200;
	add.s32 	%r205, %r204, %r186;
	and.b32  	%r206, %r204, 31;
	shf.l.wrap.b32 	%r207, %r205, %r205, %r206;
	add.s32 	%r208, %r203, %r207;
	add.s32 	%r209, %r208, -1089828067;
	shf.l.wrap.b32 	%r210, %r209, %r209, 3;
	add.s32 	%r211, %r210, %r207;
	add.s32 	%r212, %r211, %r193;
	and.b32  	%r213, %r211, 31;
	shf.l.wrap.b32 	%r214, %r212, %r212, %r213;
	add.s32 	%r215, %r210, %r678;
	add.s32 	%r216, %r215, %r214;
	shf.l.wrap.b32 	%r217, %r216, %r216, 3;
	add.s32 	%r218, %r217, %r214;
	add.s32 	%r219, %r218, %r200;
	and.b32  	%r220, %r218, 31;
	shf.l.wrap.b32 	%r221, %r219, %r219, %r220;
	add.s32 	%r222, %r217, %r41;
	add.s32 	%r223, %r222, %r221;
	shf.l.wrap.b32 	%r224, %r223, %r223, 3;
	add.s32 	%r225, %r224, %r221;
	add.s32 	%r226, %r225, %r207;
	and.b32  	%r227, %r225, 31;
	shf.l.wrap.b32 	%r228, %r226, %r226, %r227;
	add.s32 	%r229, %r224, %r49;
	add.s32 	%r230, %r229, %r228;
	shf.l.wrap.b32 	%r231, %r230, %r230, 3;
	add.s32 	%r232, %r231, %r228;
	add.s32 	%r233, %r232, %r214;
	and.b32  	%r234, %r232, 31;
	shf.l.wrap.b32 	%r235, %r233, %r233, %r234;
	add.s32 	%r236, %r231, %r56;
	add.s32 	%r237, %r236, %r235;
	shf.l.wrap.b32 	%r238, %r237, %r237, 3;
	add.s32 	%r239, %r238, %r235;
	add.s32 	%r240, %r239, %r221;
	and.b32  	%r241, %r239, 31;
	shf.l.wrap.b32 	%r242, %r240, %r240, %r241;
	add.s32 	%r243, %r238, %r63;
	add.s32 	%r244, %r243, %r242;
	shf.l.wrap.b32 	%r245, %r244, %r244, 3;
	add.s32 	%r246, %r245, %r242;
	add.s32 	%r247, %r246, %r228;
	and.b32  	%r248, %r246, 31;
	shf.l.wrap.b32 	%r249, %r247, %r247, %r248;
	add.s32 	%r250, %r245, %r70;
	add.s32 	%r251, %r250, %r249;
	shf.l.wrap.b32 	%r252, %r251, %r251, 3;
	add.s32 	%r253, %r252, %r249;
	add.s32 	%r254, %r253, %r235;
	and.b32  	%r255, %r253, 31;
	shf.l.wrap.b32 	%r256, %r254, %r254, %r255;
	add.s32 	%r257, %r252, %r77;
	add.s32 	%r258, %r257, %r256;
	shf.l.wrap.b32 	%r259, %r258, %r258, 3;
	add.s32 	%r260, %r259, %r256;
	add.s32 	%r261, %r260, %r242;
	and.b32  	%r262, %r260, 31;
	shf.l.wrap.b32 	%r263, %r261, %r261, %r262;
	add.s32 	%r264, %r259, %r84;
	add.s32 	%r265, %r264, %r263;
	shf.l.wrap.b32 	%r266, %r265, %r265, 3;
	add.s32 	%r267, %r266, %r263;
	add.s32 	%r268, %r267, %r249;
	and.b32  	%r269, %r267, 31;
	shf.l.wrap.b32 	%r270, %r268, %r268, %r269;
	add.s32 	%r271, %r266, %r91;
	add.s32 	%r272, %r271, %r270;
	shf.l.wrap.b32 	%r273, %r272, %r272, 3;
	add.s32 	%r274, %r273, %r270;
	add.s32 	%r275, %r274, %r256;
	and.b32  	%r276, %r274, 31;
	shf.l.wrap.b32 	%r277, %r275, %r275, %r276;
	add.s32 	%r278, %r273, %r98;
	add.s32 	%r279, %r278, %r277;
	shf.l.wrap.b32 	%r280, %r279, %r279, 3;
	add.s32 	%r281, %r280, %r277;
	add.s32 	%r282, %r281, %r263;
	and.b32  	%r283, %r281, 31;
	shf.l.wrap.b32 	%r284, %r282, %r282, %r283;
	add.s32 	%r285, %r280, %r105;
	add.s32 	%r286, %r285, %r284;
	shf.l.wrap.b32 	%r287, %r286, %r286, 3;
	add.s32 	%r288, %r287, %r284;
	add.s32 	%r289, %r288, %r270;
	and.b32  	%r290, %r288, 31;
	shf.l.wrap.b32 	%r291, %r289, %r289, %r290;
	add.s32 	%r292, %r287, %r112;
	add.s32 	%r293, %r292, %r291;
	shf.l.wrap.b32 	%r294, %r293, %r293, 3;
	add.s32 	%r295, %r294, %r291;
	add.s32 	%r296, %r295, %r277;
	and.b32  	%r297, %r295, 31;
	shf.l.wrap.b32 	%r298, %r296, %r296, %r297;
	add.s32 	%r299, %r294, %r119;
	add.s32 	%r300, %r299, %r298;
	shf.l.wrap.b32 	%r301, %r300, %r300, 3;
	add.s32 	%r302, %r301, %r298;
	add.s32 	%r303, %r302, %r284;
	and.b32  	%r304, %r302, 31;
	shf.l.wrap.b32 	%r305, %r303, %r303, %r304;
	add.s32 	%r306, %r301, %r126;
	add.s32 	%r307, %r306, %r305;
	shf.l.wrap.b32 	%r308, %r307, %r307, 3;
	add.s32 	%r309, %r308, %r305;
	add.s32 	%r310, %r309, %r291;
	and.b32  	%r311, %r309, 31;
	shf.l.wrap.b32 	%r312, %r310, %r310, %r311;
	add.s32 	%r313, %r308, %r133;
	add.s32 	%r314, %r313, %r312;
	shf.l.wrap.b32 	%r315, %r314, %r314, 3;
	add.s32 	%r316, %r315, %r312;
	add.s32 	%r317, %r316, %r298;
	and.b32  	%r318, %r316, 31;
	shf.l.wrap.b32 	%r319, %r317, %r317, %r318;
	add.s32 	%r320, %r315, %r140;
	add.s32 	%r321, %r320, %r319;
	shf.l.wrap.b32 	%r322, %r321, %r321, 3;
	add.s32 	%r323, %r322, %r319;
	add.s32 	%r324, %r323, %r305;
	and.b32  	%r325, %r323, 31;
	shf.l.wrap.b32 	%r326, %r324, %r324, %r325;
	add.s32 	%r327, %r322, %r147;
	add.s32 	%r328, %r327, %r326;
	shf.l.wrap.b32 	%r329, %r328, %r328, 3;
	add.s32 	%r330, %r329, %r326;
	add.s32 	%r331, %r330, %r312;
	and.b32  	%r332, %r330, 31;
	shf.l.wrap.b32 	%r333, %r331, %r331, %r332;
	add.s32 	%r334, %r329, %r154;
	add.s32 	%r335, %r334, %r333;
	shf.l.wrap.b32 	%r336, %r335, %r335, 3;
	add.s32 	%r337, %r336, %r333;
	add.s32 	%r338, %r337, %r319;
	and.b32  	%r339, %r337, 31;
	shf.l.wrap.b32 	%r340, %r338, %r338, %r339;
	add.s32 	%r341, %r336, %r161;
	add.s32 	%r342, %r341, %r340;
	shf.l.wrap.b32 	%r343, %r342, %r342, 3;
	add.s32 	%r344, %r343, %r340;
	add.s32 	%r345, %r344, %r326;
	and.b32  	%r346, %r344, 31;
	shf.l.wrap.b32 	%r347, %r345, %r345, %r346;
	add.s32 	%r348, %r343, %r168;
	add.s32 	%r349, %r348, %r347;
	shf.l.wrap.b32 	%r350, %r349, %r349, 3;
	add.s32 	%r351, %r350, %r347;
	add.s32 	%r352, %r351, %r333;
	and.b32  	%r353, %r351, 31;
	shf.l.wrap.b32 	%r354, %r352, %r352, %r353;
	add.s32 	%r355, %r350, %r175;
	add.s32 	%r356, %r355, %r354;
	shf.l.wrap.b32 	%r357, %r356, %r356, 3;
	add.s32 	%r358, %r357, %r354;
	add.s32 	%r359, %r358, %r340;
	and.b32  	%r360, %r358, 31;
	shf.l.wrap.b32 	%r361, %r359, %r359, %r360;
	add.s32 	%r362, %r357, %r182;
	add.s32 	%r363, %r362, %r361;
	shf.l.wrap.b32 	%r364, %r363, %r363, 3;
	add.s32 	%r365, %r364, %r361;
	add.s32 	%r366, %r365, %r347;
	and.b32  	%r367, %r365, 31;
	shf.l.wrap.b32 	%r368, %r366, %r366, %r367;
	add.s32 	%r369, %r364, %r189;
	add.s32 	%r370, %r369, %r368;
	shf.l.wrap.b32 	%r371, %r370, %r370, 3;
	add.s32 	%r372, %r371, %r368;
	add.s32 	%r373, %r372, %r354;
	and.b32  	%r374, %r372, 31;
	shf.l.wrap.b32 	%r375, %r373, %r373, %r374;
	add.s32 	%r376, %r371, %r196;
	add.s32 	%r377, %r376, %r375;
	shf.l.wrap.b32 	%r378, %r377, %r377, 3;
	add.s32 	%r379, %r378, %r375;
	add.s32 	%r380, %r379, %r361;
	and.b32  	%r381, %r379, 31;
	shf.l.wrap.b32 	%r382, %r380, %r380, %r381;
	add.s32 	%r383, %r378, %r203;
	add.s32 	%r384, %r383, %r382;
	shf.l.wrap.b32 	%r10, %r384, %r384, 3;
	add.s32 	%r385, %r10, %r382;
	add.s32 	%r386, %r385, %r368;
	and.b32  	%r387, %r385, 31;
	shf.l.wrap.b32 	%r388, %r386, %r386, %r387;
	add.s32 	%r389, %r10, %r210;
	add.s32 	%r390, %r389, %r388;
	shf.l.wrap.b32 	%r391, %r390, %r390, 3;
	add.s32 	%r392, %r391, %r388;
	add.s32 	%r393, %r392, %r375;
	and.b32  	%r394, %r392, 31;
	shf.l.wrap.b32 	%r395, %r393, %r393, %r394;
	add.s32 	%r396, %r391, %r217;
	add.s32 	%r397, %r396, %r395;
	shf.l.wrap.b32 	%r398, %r397, %r397, 3;
	add.s32 	%r399, %r398, %r395;
	add.s32 	%r400, %r399, %r382;
	and.b32  	%r401, %r399, 31;
	shf.l.wrap.b32 	%r402, %r400, %r400, %r401;
	add.s32 	%r403, %r398, %r224;
	add.s32 	%r404, %r403, %r402;
	shf.l.wrap.b32 	%r405, %r404, %r404, 3;
	add.s32 	%r406, %r405, %r402;
	add.s32 	%r407, %r406, %r388;
	and.b32  	%r408, %r406, 31;
	shf.l.wrap.b32 	%r409, %r407, %r407, %r408;
	add.s32 	%r410, %r405, %r231;
	add.s32 	%r411, %r410, %r409;
	shf.l.wrap.b32 	%r412, %r411, %r411, 3;
	add.s32 	%r413, %r412, %r409;
	add.s32 	%r414, %r413, %r395;
	and.b32  	%r415, %r413, 31;
	shf.l.wrap.b32 	%r416, %r414, %r414, %r415;
	add.s32 	%r417, %r412, %r238;
	add.s32 	%r418, %r417, %r416;
	shf.l.wrap.b32 	%r419, %r418, %r418, 3;
	add.s32 	%r420, %r419, %r416;
	add.s32 	%r421, %r420, %r402;
	and.b32  	%r422, %r420, 31;
	shf.l.wrap.b32 	%r423, %r421, %r421, %r422;
	add.s32 	%r424, %r419, %r245;
	add.s32 	%r425, %r424, %r423;
	shf.l.wrap.b32 	%r426, %r425, %r425, 3;
	add.s32 	%r427, %r426, %r423;
	add.s32 	%r428, %r427, %r409;
	and.b32  	%r429, %r427, 31;
	shf.l.wrap.b32 	%r430, %r428, %r428, %r429;
	add.s32 	%r431, %r426, %r252;
	add.s32 	%r432, %r431, %r430;
	shf.l.wrap.b32 	%r433, %r432, %r432, 3;
	add.s32 	%r434, %r433, %r430;
	add.s32 	%r435, %r434, %r416;
	and.b32  	%r436, %r434, 31;
	shf.l.wrap.b32 	%r437, %r435, %r435, %r436;
	add.s32 	%r438, %r433, %r259;
	add.s32 	%r439, %r438, %r437;
	shf.l.wrap.b32 	%r440, %r439, %r439, 3;
	add.s32 	%r441, %r440, %r437;
	add.s32 	%r442, %r441, %r423;
	and.b32  	%r443, %r441, 31;
	shf.l.wrap.b32 	%r444, %r442, %r442, %r443;
	add.s32 	%r445, %r440, %r266;
	add.s32 	%r446, %r445, %r444;
	shf.l.wrap.b32 	%r447, %r446, %r446, 3;
	add.s32 	%r448, %r447, %r444;
	add.s32 	%r449, %r448, %r430;
	and.b32  	%r450, %r448, 31;
	shf.l.wrap.b32 	%r451, %r449, %r449, %r450;
	add.s32 	%r452, %r447, %r273;
	add.s32 	%r453, %r452, %r451;
	shf.l.wrap.b32 	%r454, %r453, %r453, 3;
	add.s32 	%r455, %r454, %r451;
	add.s32 	%r456, %r455, %r437;
	and.b32  	%r457, %r455, 31;
	shf.l.wrap.b32 	%r458, %r456, %r456, %r457;
	add.s32 	%r459, %r454, %r280;
	add.s32 	%r460, %r459, %r458;
	shf.l.wrap.b32 	%r461, %r460, %r460, 3;
	add.s32 	%r462, %r461, %r458;
	add.s32 	%r463, %r462, %r444;
	and.b32  	%r464, %r462, 31;
	shf.l.wrap.b32 	%r465, %r463, %r463, %r464;
	add.s32 	%r466, %r461, %r287;
	add.s32 	%r467, %r466, %r465;
	shf.l.wrap.b32 	%r468, %r467, %r467, 3;
	add.s32 	%r469, %r468, %r465;
	add.s32 	%r470, %r469, %r451;
	and.b32  	%r471, %r469, 31;
	shf.l.wrap.b32 	%r472, %r470, %r470, %r471;
	add.s32 	%r473, %r468, %r294;
	add.s32 	%r474, %r473, %r472;
	shf.l.wrap.b32 	%r475, %r474, %r474, 3;
	add.s32 	%r476, %r475, %r472;
	add.s32 	%r477, %r476, %r458;
	and.b32  	%r478, %r476, 31;
	shf.l.wrap.b32 	%r479, %r477, %r477, %r478;
	add.s32 	%r480, %r475, %r301;
	add.s32 	%r481, %r480, %r479;
	shf.l.wrap.b32 	%r482, %r481, %r481, 3;
	add.s32 	%r483, %r482, %r479;
	add.s32 	%r484, %r483, %r465;
	and.b32  	%r485, %r483, 31;
	shf.l.wrap.b32 	%r486, %r484, %r484, %r485;
	add.s32 	%r487, %r482, %r308;
	add.s32 	%r488, %r487, %r486;
	shf.l.wrap.b32 	%r489, %r488, %r488, 3;
	add.s32 	%r490, %r489, %r486;
	add.s32 	%r491, %r490, %r472;
	and.b32  	%r492, %r490, 31;
	shf.l.wrap.b32 	%r493, %r491, %r491, %r492;
	add.s32 	%r494, %r489, %r315;
	add.s32 	%r495, %r494, %r493;
	shf.l.wrap.b32 	%r496, %r495, %r495, 3;
	add.s32 	%r497, %r496, %r493;
	add.s32 	%r498, %r497, %r479;
	and.b32  	%r499, %r497, 31;
	shf.l.wrap.b32 	%r500, %r498, %r498, %r499;
	add.s32 	%r501, %r496, %r322;
	add.s32 	%r502, %r501, %r500;
	shf.l.wrap.b32 	%r503, %r502, %r502, 3;
	add.s32 	%r504, %r503, %r500;
	add.s32 	%r505, %r504, %r486;
	and.b32  	%r506, %r504, 31;
	shf.l.wrap.b32 	%r507, %r505, %r505, %r506;
	add.s32 	%r508, %r503, %r329;
	add.s32 	%r509, %r508, %r507;
	shf.l.wrap.b32 	%r510, %r509, %r509, 3;
	add.s32 	%r511, %r510, %r507;
	add.s32 	%r512, %r511, %r493;
	and.b32  	%r513, %r511, 31;
	shf.l.wrap.b32 	%r514, %r512, %r512, %r513;
	add.s32 	%r515, %r510, %r336;
	add.s32 	%r516, %r515, %r514;
	shf.l.wrap.b32 	%r517, %r516, %r516, 3;
	add.s32 	%r518, %r517, %r514;
	add.s32 	%r519, %r518, %r500;
	and.b32  	%r520, %r518, 31;
	shf.l.wrap.b32 	%r521, %r519, %r519, %r520;
	add.s32 	%r522, %r517, %r343;
	add.s32 	%r523, %r522, %r521;
	shf.l.wrap.b32 	%r524, %r523, %r523, 3;
	add.s32 	%r525, %r524, %r521;
	add.s32 	%r526, %r525, %r507;
	and.b32  	%r527, %r525, 31;
	shf.l.wrap.b32 	%r528, %r526, %r526, %r527;
	add.s32 	%r529, %r524, %r350;
	add.s32 	%r530, %r529, %r528;
	shf.l.wrap.b32 	%r531, %r530, %r530, 3;
	add.s32 	%r532, %r531, %r528;
	add.s32 	%r533, %r532, %r514;
	and.b32  	%r534, %r532, 31;
	shf.l.wrap.b32 	%r535, %r533, %r533, %r534;
	add.s32 	%r536, %r531, %r357;
	add.s32 	%r537, %r536, %r535;
	shf.l.wrap.b32 	%r538, %r537, %r537, 3;
	add.s32 	%r539, %r538, %r535;
	add.s32 	%r540, %r539, %r521;
	and.b32  	%r541, %r539, 31;
	shf.l.wrap.b32 	%r11, %r540, %r540, %r541;
	add.s32 	%r542, %r538, %r364;
	add.s32 	%r543, %r542, %r11;
	shf.l.wrap.b32 	%r544, %r543, %r543, 3;
	add.s32 	%r545, %r544, %r11;
	add.s32 	%r546, %r545, %r528;
	and.b32  	%r547, %r545, 31;
	shf.l.wrap.b32 	%r548, %r546, %r546, %r547;
	add.s32 	%r549, %r544, %r371;
	add.s32 	%r550, %r549, %r548;
	shf.l.wrap.b32 	%r551, %r550, %r550, 3;
	add.s32 	%r552, %r551, %r548;
	add.s32 	%r553, %r552, %r535;
	and.b32  	%r554, %r552, 31;
	shf.l.wrap.b32 	%r12, %r553, %r553, %r554;
	add.s32 	%r555, %r551, %r378;
	add.s32 	%r556, %r555, %r12;
	shf.l.wrap.b32 	%r13, %r556, %r556, 3;
	ld.const.u32 	%r557, [%rd2+16];
	add.s32 	%r558, %r391, %r557;
	ld.const.u32 	%r559, [%rd2+20];
	add.s32 	%r560, %r398, %r559;
	xor.b32  	%r561, %r560, %r558;
	and.b32  	%r562, %r560, 31;
	shf.l.wrap.b32 	%r563, %r561, %r561, %r562;
	add.s32 	%r564, %r405, %r563;
	xor.b32  	%r565, %r564, %r560;
	and.b32  	%r566, %r564, 31;
	shf.l.wrap.b32 	%r567, %r565, %r565, %r566;
	add.s32 	%r568, %r412, %r567;
	xor.b32  	%r569, %r568, %r564;
	and.b32  	%r570, %r568, 31;
	shf.l.wrap.b32 	%r571, %r569, %r569, %r570;
	add.s32 	%r572, %r419, %r571;
	xor.b32  	%r573, %r572, %r568;
	and.b32  	%r574, %r572, 31;
	shf.l.wrap.b32 	%r575, %r573, %r573, %r574;
	add.s32 	%r576, %r426, %r575;
	xor.b32  	%r577, %r576, %r572;
	and.b32  	%r578, %r576, 31;
	shf.l.wrap.b32 	%r579, %r577, %r577, %r578;
	add.s32 	%r580, %r433, %r579;
	xor.b32  	%r581, %r580, %r576;
	and.b32  	%r582, %r580, 31;
	shf.l.wrap.b32 	%r583, %r581, %r581, %r582;
	add.s32 	%r584, %r440, %r583;
	xor.b32  	%r585, %r584, %r580;
	and.b32  	%r586, %r584, 31;
	shf.l.wrap.b32 	%r587, %r585, %r585, %r586;
	add.s32 	%r588, %r447, %r587;
	xor.b32  	%r589, %r588, %r584;
	and.b32  	%r590, %r588, 31;
	shf.l.wrap.b32 	%r591, %r589, %r589, %r590;
	add.s32 	%r592, %r454, %r591;
	xor.b32  	%r593, %r592, %r588;
	and.b32  	%r594, %r592, 31;
	shf.l.wrap.b32 	%r595, %r593, %r593, %r594;
	add.s32 	%r596, %r461, %r595;
	xor.b32  	%r597, %r596, %r592;
	and.b32  	%r598, %r596, 31;
	shf.l.wrap.b32 	%r599, %r597, %r597, %r598;
	add.s32 	%r600, %r468, %r599;
	xor.b32  	%r601, %r600, %r596;
	and.b32  	%r602, %r600, 31;
	shf.l.wrap.b32 	%r603, %r601, %r601, %r602;
	add.s32 	%r604, %r475, %r603;
	xor.b32  	%r605, %r604, %r600;
	and.b32  	%r606, %r604, 31;
	shf.l.wrap.b32 	%r607, %r605, %r605, %r606;
	add.s32 	%r608, %r482, %r607;
	xor.b32  	%r609, %r608, %r604;
	and.b32  	%r610, %r608, 31;
	shf.l.wrap.b32 	%r611, %r609, %r609, %r610;
	add.s32 	%r612, %r489, %r611;
	xor.b32  	%r613, %r612, %r608;
	and.b32  	%r614, %r612, 31;
	shf.l.wrap.b32 	%r615, %r613, %r613, %r614;
	add.s32 	%r616, %r496, %r615;
	xor.b32  	%r617, %r616, %r612;
	and.b32  	%r618, %r616, 31;
	shf.l.wrap.b32 	%r619, %r617, %r617, %r618;
	add.s32 	%r620, %r503, %r619;
	xor.b32  	%r621, %r620, %r616;
	and.b32  	%r622, %r620, 31;
	shf.l.wrap.b32 	%r623, %r621, %r621, %r622;
	add.s32 	%r624, %r510, %r623;
	xor.b32  	%r625, %r624, %r620;
	and.b32  	%r626, %r624, 31;
	shf.l.wrap.b32 	%r627, %r625, %r625, %r626;
	add.s32 	%r628, %r517, %r627;
	xor.b32  	%r629, %r628, %r624;
	and.b32  	%r630, %r628, 31;
	shf.l.wrap.b32 	%r631, %r629, %r629, %r630;
	add.s32 	%r632, %r524, %r631;
	xor.b32  	%r633, %r632, %r628;
	and.b32  	%r634, %r632, 31;
	shf.l.wrap.b32 	%r635, %r633, %r633, %r634;
	add.s32 	%r636, %r531, %r635;
	xor.b32  	%r637, %r636, %r632;
	and.b32  	%r638, %r636, 31;
	shf.l.wrap.b32 	%r639, %r637, %r637, %r638;
	add.s32 	%r640, %r538, %r639;
	xor.b32  	%r641, %r640, %r636;
	and.b32  	%r642, %r640, 31;
	shf.l.wrap.b32 	%r643, %r641, %r641, %r642;
	add.s32 	%r644, %r544, %r643;
	xor.b32  	%r645, %r644, %r640;
	and.b32  	%r646, %r644, 31;
	shf.l.wrap.b32 	%r647, %r645, %r645, %r646;
	add.s32 	%r14, %r551, %r647;
	xor.b32  	%r648, %r14, %r644;
	and.b32  	%r649, %r14, 31;
	shf.l.wrap.b32 	%r650, %r648, %r648, %r649;
	add.s32 	%r15, %r13, %r650;
	ld.const.u32 	%r651, [%rd2+24];
	setp.ne.s32 	%p2, %r15, %r651;
	@%p2 bra 	$L__BB0_4;

	mov.b32 	%r676, %envreg3;
	mov.u32 	%r675, %tid.x;
	add.s32 	%r674, %r675, %r676;
	mov.u32 	%r673, %ctaid.x;
	mov.u32 	%r672, %ntid.x;
	mad.lo.s32 	%r671, %r672, %r673, %r674;
	ld.param.u64 	%rd9, [ocl_rc572_1pipe_param_1];
	ld.param.u64 	%rd8, [ocl_rc572_1pipe_param_0];
	add.s32 	%r652, %r13, %r12;
	add.s32 	%r653, %r652, %r11;
	and.b32  	%r654, %r652, 31;
	shf.l.wrap.b32 	%r655, %r653, %r653, %r654;
	add.s32 	%r656, %r13, %r10;
	add.s32 	%r657, %r656, %r655;
	shf.l.wrap.b32 	%r658, %r657, %r657, 3;
	and.b32  	%r659, %r15, 31;
	xor.b32  	%r660, %r15, %r14;
	shf.l.wrap.b32 	%r661, %r660, %r660, %r659;
	add.s32 	%r662, %r658, %r661;
	atom.global.add.u32 	%r663, [%rd9], 1;
	shl.b32 	%r664, %r663, 1;
	or.b32  	%r665, %r664, 1;
	ld.const.u32 	%r666, [%rd8+12];
	add.s32 	%r667, %r666, %r671;
	ld.const.u32 	%r668, [%rd8+28];
	setp.eq.s32 	%p3, %r662, %r668;
	selp.b32 	%r669, -2147483648, 0, %p3;
	mul.wide.u32 	%rd4, %r665, 4;
	add.s64 	%rd5, %rd9, %rd4;
	st.global.u32 	[%rd5], %r669;
	add.s32 	%r670, %r664, 2;
	mul.wide.u32 	%rd6, %r670, 4;
	add.s64 	%rd7, %rd9, %rd6;
	st.global.u32 	[%rd7], %r667;

$L__BB0_4:
	ret;

}

\00\00\00
6 
	.headerflags	@"EF_CUDA_TEXMODE_INDEPENDENT EF_CUDA_64BIT_ADDRESS EF_CUDA_SM75 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM75)"
	.elftype	@"ET_EXEC"


//--------------------- .debug_frame              --------------------------
	.section	.debug_frame,"",@progbits
.debug_frame:
        /*0000*/ 	.byte	0xff, 0xff, 0xff, 0xff, 0x24, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff
        /*0010*/ 	.byte	0xff, 0xff, 0xff, 0xff, 0x03, 0x00, 0x04, 0x7c, 0xff, 0xff, 0xff, 0xff, 0x0f, 0x0c, 0x81, 0x80
        /*0020*/ 	.byte	0x80, 0x28, 0x00, 0x08, 0xff, 0x81, 0x80, 0x28, 0x08, 0x81, 0x80, 0x80, 0x28, 0x00, 0x00, 0x00
        /*0030*/ 	.byte	0xff, 0xff, 0xff, 0xff, 0x34, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00
        /*0040*/ 	.byte	0x00, 0x00, 0x00, 0x00
        /*0044*/ 	.dword	ocl_rc572_1pipe
        /*004c*/ 	.byte	0x80, 0x26, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x04, 0x00, 0x00, 0x00, 0x04, 0xf0, 0x08
        /*005c*/ 	.byte	0x00, 0x00, 0x0c, 0x81, 0x80, 0x80, 0x28, 0x00, 0x04, 0x88, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00
        /*006c*/ 	.byte	0x00, 0x00, 0x00, 0x00


//--------------------- .nv.info                  --------------------------
	.section	.nv.info,"",@"SHT_CUDA_INFO"
	.align	4


	//----- nvinfo : EIATTR_REGCOUNT
	.align		4
        /*0000*/ 	.byte	0x04, 0x2f
        /*0002*/ 	.short	(.L_1 - .L_0)
	.align		4
.L_0:
        /*0004*/ 	.word	index@(ocl_rc572_1pipe)
        /*0008*/ 	.word	0x00000026


	//----- nvinfo : EIATTR_MAX_STACK_SIZE
	.align		4
.L_1:
        /*000c*/ 	.byte	0x04, 0x23
        /*000e*/ 	.short	(.L_3 - .L_2)
	.align		4
.L_2:
        /*0010*/ 	.word	index@(ocl_rc572_1pipe)
        /*0014*/ 	.word	0x00000000


	//----- nvinfo : EIATTR_MIN_STACK_SIZE
	.align		4
.L_3:
        /*0018*/ 	.byte	0x04, 0x12
        /*001a*/ 	.short	(.L_5 - .L_4)
	.align		4
.L_4:
        /*001c*/ 	.word	index@(ocl_rc572_1pipe)
        /*0020*/ 	.word	0x00000000


	//----- nvinfo : EIATTR_FRAME_SIZE
	.align		4
.L_5:
        /*0024*/ 	.byte	0x04, 0x11
        /*0026*/ 	.short	(.L_7 - .L_6)
	.align		4
.L_6:
        /*0028*/ 	.word	index@(ocl_rc572_1pipe)
        /*002c*/ 	.word	0x00000000
.L_7:


//--------------------- .nv.info.ocl_rc572_1pipe  --------------------------
	.section	.nv.info.ocl_rc572_1pipe,"",@"SHT_CUDA_INFO"
	.align	4


	//----- nvinfo : EIATTR_SW_WAR
	.align		4
        /*0000*/ 	.byte	0x04, 0x36
        /*0002*/ 	.short	(.L_9 - .L_8)
.L_8:
        /*0004*/ 	.word	0x00000001


	//----- nvinfo : EIATTR_CUDA_API_VERSION
	.align		4
.L_9:
        /*0008*/ 	.byte	0x04, 0x37
        /*000a*/ 	.short	(.L_11 - .L_10)
.L_10:
        /*000c*/ 	.word	0x00000078


	//----- nvinfo : EIATTR_PARAM_CBANK
	.align		4
.L_11:
        /*0010*/ 	.byte	0x04, 0x0a
        /*0012*/ 	.short	(.L_13 - .L_12)
	.align		4
.L_12:
        /*0014*/ 	.word	index@(.nv.constant0.ocl_rc572_1pipe)
        /*0018*/ 	.short	0x0160
        /*001a*/ 	.short	0x0010


	//----- nvinfo : EIATTR_CBANK_PARAM_SIZE
	.align		4
.L_13:
        /*001c*/ 	.byte	0x03, 0x19
        /*001e*/ 	.short	0x0010


	//----- nvinfo : EIATTR_KPARAM_INFO
	.align		4
        /*0020*/ 	.byte	0x04, 0x17
        /*0022*/ 	.short	(.L_15 - .L_14)
.L_14:
        /*0024*/ 	.word	0x00000000
        /*0028*/ 	.short	0x0001
        /*002a*/ 	.short	0x0008
        /*002c*/ 	.byte	0x02, 0xf4, 0x21, 0x00


	//----- nvinfo : EIATTR_KPARAM_INFO
	.align		4
.L_15:
        /*0030*/ 	.byte	0x04, 0x17
        /*0032*/ 	.short	(.L_17 - .L_16)
.L_16:
        /*0034*/ 	.word	0x00000000
        /*0038*/ 	.short	0x0000
        /*003a*/ 	.short	0x0000
        /*003c*/ 	.byte	0x02, 0xf3, 0x21, 0x00


	//----- nvinfo : EIATTR_MAXREG_COUNT
	.align		4
.L_17:
        /*0040*/ 	.byte	0x03, 0x1b
        /*0042*/ 	.short	0x00ff


	//----- nvinfo : EIATTR_INT_WARP_WIDE_INSTR_OFFSETS
	.align		4
        /*0044*/ 	.byte	0x04, 0x31
        /*0046*/ 	.short	(.L_19 - .L_18)


	//   ....[0]....
.L_18:
        /*0048*/ 	.word	0x000023e0


	//   ....[1]....
        /*004c*/ 	.word	0x00002550


	//----- nvinfo : EIATTR_EXIT_INSTR_OFFSETS
	.align		4
.L_19:
        /*0050*/ 	.byte	0x04, 0x1c
        /*0052*/ 	.short	(.L_21 - .L_20)


	//   ....[0]....
.L_20:
        /*0054*/ 	.word	0x000023c0


	//   ....[1]....
        /*0058*/ 	.word	0x000025f0
.L_21:


//--------------------- .nv.constant0.ocl_rc572_1pipe --------------------------
	.section	.nv.constant0.ocl_rc572_1pipe,"a",@progbits
	.align	4
.nv.constant0.ocl_rc572_1pipe:
	.zero		368


//--------------------- .text.ocl_rc572_1pipe     --------------------------
	.section	.text.ocl_rc572_1pipe,"ax",@progbits
	.sectioninfo	@"SHI_REGISTERS=38"
	.align	128
.text.ocl_rc572_1pipe:
        .type           ocl_rc572_1pipe,@function
        .size           ocl_rc572_1pipe,(.L_x_1 - ocl_rc572_1pipe)
        .other          ocl_rc572_1pipe,@"STO_CUDA_ENTRY STV_DEFAULT"
ocl_rc572_1pipe:
        /*0000*/                   IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
        /*0010*/                   S2R R2, SR_TID.X ;
        /*0020*/                   IMAD.MOV.U32 R0, RZ, RZ, c[0x0][0x160] ;
        /*0030*/                   S2R R5, SR_CTAID.X ;
        /*0040*/                   LDC.IL R3, c[0x0][R0] ;
        /*0050*/                   LDC.IL R4, c[0x0][R0+0x4] ;
        /*0060*/                   LDC.IL R6, c[0x0][R0+0x20] ;
        /*0070*/                   IADD3 R2, R2, c[0x0][0x94], RZ ;
        /*0080*/                   IMAD R2, R5, c[0x0][0x0], R2 ;
        /*0090*/                   LDC.IL R35, c[0x0][R0+0x10] ;
        /*00a0*/                   IMAD.IADD R5, R3, 0x1, R2 ;
        /*00b0*/                   SHF.R.U32.HI R3, RZ, 0x8, R5 ;
        /*00c0*/                   LDC.IL R34, c[0x0][R0+0x14] ;
        /*00d0*/                   LOP3.LUT R8, R5, 0xff, RZ, 0xc0, !PT ;
        /*00e0*/                   IMAD.IADD R4, R3, 0x1, R4 ;
        /*00f0*/                   ISETP.GE.U32.AND P0, PT, R4, R3, PT ;
        /*0100*/                   LDC.IL R3, c[0x0][R0+0x24] ;
        /*0110*/              @!P0 LDC.IL R7, c[0x0][R0+0x8] ;
        /*0120*/              @!P0 PRMT R7, R7, 0x123, RZ ;
        /*0130*/              @!P0 IADD3 R7, R7, 0x1, RZ ;
        /*0140*/              @!P0 PRMT R7, R7, 0x123, RZ ;
        /*0150*/              @!P0 IADD3 R7, R7, -0x40f574e3, RZ ;
        /*0160*/              @!P0 SHF.L.W.U32.HI R6, R7, 0x1d, R7 ;
        /*0170*/              @!P0 IADD3 R7, R6, 0x15235639, RZ ;
        /*0180*/              @!P0 SHF.L.W.U32.HI R3, R7, 0x3, R7 ;
        /*0190*/                   PRMT R7, R4, 0x123, RZ ;
        /*01a0*/                   IMAD.IADD R4, R3, 0x1, R6 ;
        /*01b0*/                   IMAD.IADD R7, R4.reuse, 0x1, R7 ;
        /*01c0*/                   LOP3.LUT R4, R4, 0x1f, RZ, 0xc0, !PT ;
        /*01d0*/                   SHF.L.W.U32.HI R7, R7, R4, R7 ;
        /*01e0*/                   IADD3 R4, R3, -0xbafbb2b, R7 ;
        /*01f0*/                   SHF.L.W.U32.HI R4, R4, 0x3, R4 ;
        /*0200*/                   IMAD.IADD R5, R7, 0x1, R4 ;
        /*0210*/                   IMAD.IADD R8, R5.reuse, 0x1, R8 ;
        /*0220*/                   LOP3.LUT R5, R5, 0x1f, RZ, 0xc0, !PT ;
        /*0230*/                   SHF.L.W.U32.HI R8, R8, R5, R8 ;
        /*0240*/                   IADD3 R5, R4, -0x6d784172, R8 ;
        /*0250*/                   SHF.L.W.U32.HI R5, R5, 0x3, R5 ;
        /*0260*/                   IMAD.IADD R9, R8, 0x1, R5 ;
        /*0270*/                   IMAD.IADD R6, R9.reuse, 0x1, R6 ;
        /*0280*/                   LOP3.LUT R9, R9, 0x1f, RZ, 0xc0, !PT ;
        /*0290*/                   SHF.L.W.U32.HI R9, R6, R9, R6 ;
        /*02a0*/                   IADD3 R6, R5, 0x30bf3847, R9 ;
        /*02b0*/                   SHF.L.W.U32.HI R6, R6, 0x3, R6 ;
        /*02c0*/                   IMAD.IADD R10, R9, 0x1, R6 ;
        /*02d0*/                   IMAD.IADD R7, R7, 0x1, R10 ;
        /*02e0*/                   LOP3.LUT R10, R10, 0x1f, RZ, 0xc0, !PT ;
        /*02f0*/                   SHF.L.W.U32.HI R10, R7, R10, R7 ;
        /*0300*/                   IADD3 R7, R6, -0x31094e00, R10 ;
        /*0310*/                   SHF.L.W.U32.HI R7, R7, 0x3, R7 ;
        /*0320*/                   IMAD.IADD R11, R10, 0x1, R7 ;
        /*0330*/                   IMAD.IADD R8, R8, 0x1, R11 ;
        /*0340*/                   LOP3.LUT R11, R11, 0x1f, RZ, 0xc0, !PT ;
        /*0350*/                   SHF.L.W.U32.HI R11, R8, R11, R8 ;
        /*0360*/                   IADD3 R8, R7, 0x6d2e2bb9, R11 ;
        /*0370*/                   SHF.L.W.U32.HI R8, R8, 0x3, R8 ;
        /*0380*/                   IMAD.IADD R12, R11, 0x1, R8 ;
        /*0390*/                   IMAD.IADD R9, R9, 0x1, R12 ;
        /*03a0*/                   LOP3.LUT R12, R12, 0x1f, RZ, 0xc0, !PT ;
        /*03b0*/                   SHF.L.W.U32.HI R12, R9, R12, R9 ;
        /*03c0*/                   IADD3 R9, R8, 0xb65a572, R12 ;
        /*03d0*/                   SHF.L.W.U32.HI R9, R9, 0x3, R9 ;
        /*03e0*/                   IMAD.IADD R13, R12, 0x1, R9 ;
        /*03f0*/                   IMAD.IADD R10, R10, 0x1, R13 ;
        /*0400*/                   LOP3.LUT R13, R13, 0x1f, RZ, 0xc0, !PT ;
        /*0410*/                   SHF.L.W.U32.HI R13, R10, R13, R10 ;
        /*0420*/                   IADD3 R10, R9, -0x5662e0d5, R13 ;
        /*0430*/                   SHF.L.W.U32.HI R10, R10, 0x3, R10 ;
        /*0440*/                   IMAD.IADD R14, R13, 0x1, R10 ;
        /*0450*/                   IMAD.IADD R11, R11, 0x1, R14 ;
        /*0460*/                   LOP3.LUT R14, R14, 0x1f, RZ, 0xc0, !PT ;
        /*0470*/                   SHF.L.W.U32.HI R14, R11, R14, R11 ;
        /*0480*/                   IADD3 R11, R10, 0x47d498e4, R14 ;
        /*0490*/                   SHF.L.W.U32.HI R11, R11, 0x3, R11 ;
        /*04a0*/                   IMAD.IADD R15, R14, 0x1, R11 ;
        /*04b0*/                   IMAD.IADD R12, R12, 0x1, R15 ;
        /*04c0*/                   LOP3.LUT R15, R15, 0x1f, RZ, 0xc0, !PT ;
        /*04d0*/                   SHF.L.W.U32.HI R15, R12, R15, R12 ;
        /*04e0*/                   IADD3 R12, R11, -0x19f3ed63, R15 ;
        /*04f0*/                   SHF.L.W.U32.HI R12, R12, 0x3, R12 ;
        /*0500*/                   IMAD.IADD R16, R15, 0x1, R12 ;
        /*0510*/                   IMAD.IADD R13, R13, 0x1, R16 ;
        /*0520*/                   LOP3.LUT R16, R16, 0x1f, RZ, 0xc0, !PT ;
        /*0530*/                   SHF.L.W.U32.HI R16, R13, R16, R13 ;
        /*0540*/                   IADD3 R13, R12, -0x7bbc73aa, R16 ;
        /*0550*/                   SHF.L.W.U32.HI R13, R13, 0x3, R13 ;
        /*0560*/                   IMAD.IADD R17, R16, 0x1, R13 ;
        /*0570*/                   IMAD.IADD R14, R14, 0x1, R17 ;
        /*0580*/                   LOP3.LUT R17, R17, 0x1f, RZ, 0xc0, !PT ;
        /*0590*/                   SHF.L.W.U32.HI R17, R14, R17, R14 ;
        /*05a0*/                   IADD3 R14, R13, 0x227b060f, R17 ;
        /*05b0*/                   SHF.L.W.U32.HI R14, R14, 0x3, R14 ;
        /*05c0*/                   IMAD.IADD R18, R17, 0x1, R14 ;
        /*05d0*/                   IMAD.IADD R15, R15, 0x1, R18 ;
        /*05e0*/                   LOP3.LUT R18, R18, 0x1f, RZ, 0xc0, !PT ;
        /*05f0*/                   SHF.L.W.U32.HI R18, R15, R18, R15 ;
        /*0600*/                   IADD3 R15, R14, -0x3f4d8038, R18 ;
        /*0610*/                   SHF.L.W.U32.HI R15, R15, 0x3, R15 ;
        /*0620*/                   IMAD.IADD R19, R18, 0x1, R15 ;
        /*0630*/                   IMAD.IADD R16, R16, 0x1, R19 ;
        /*0640*/                   LOP3.LUT R19, R19, 0x1f, RZ, 0xc0, !PT ;
        /*0650*/                   SHF.L.W.U32.HI R19, R16, R19, R16 ;
        /*0660*/                   IADD3 R16, R15, 0x5ee9f981, R19 ;
        /*0670*/                   SHF.L.W.U32.HI R16, R16, 0x3, R16 ;
        /*0680*/                   IMAD.IADD R20, R19, 0x1, R16 ;
        /*0690*/                   IMAD.IADD R17, R17, 0x1, R20 ;
        /*06a0*/                   LOP3.LUT R20, R20, 0x1f, RZ, 0xc0, !PT ;
        /*06b0*/                   SHF.L.W.U32.HI R20, R17, R20, R17 ;
        /*06c0*/                   IADD3 R17, R16, -0x2de8cc6, R20 ;
        /*06d0*/                   SHF.L.W.U32.HI R17, R17, 0x3, R17 ;
        /*06e0*/                   IMAD.IADD R21, R20, 0x1, R17 ;
        /*06f0*/                   IMAD.IADD R18, R18, 0x1, R21 ;
        /*0700*/                   LOP3.LUT R21, R21, 0x1f, RZ, 0xc0, !PT ;
        /*0710*/                   SHF.L.W.U32.HI R21, R18, R21, R18 ;
        /*0720*/                   IADD3 R18, R17, -0x64a7130d, R21 ;
        /*0730*/                   SHF.L.W.U32.HI R18, R18, 0x3, R18 ;
        /*0740*/                   IMAD.IADD R22, R21, 0x1, R18 ;
        /*0750*/                   IMAD.IADD R19, R19, 0x1, R22 ;
        /*0760*/                   LOP3.LUT R22, R22, 0x1f, RZ, 0xc0, !PT ;
        /*0770*/                   SHF.L.W.U32.HI R22, R19, R22, R19 ;
        /*0780*/                   IADD3 R19, R18, 0x399066ac, R22 ;
        /*0790*/                   SHF.L.W.U32.HI R19, R19, 0x3, R19 ;
        /*07a0*/                   IMAD.IADD R23, R22, 0x1, R19 ;
        /*07b0*/                   IMAD.IADD R20, R20, 0x1, R23 ;
        /*07c0*/                   LOP3.LUT R23, R23, 0x1f, RZ, 0xc0, !PT ;
        /*07d0*/                   SHF.L.W.U32.HI R23, R20, R23, R20 ;
        /*07e0*/                   IADD3 R20, R19, -0x28381f9b, R23 ;
        /*07f0*/                   SHF.L.W.U32.HI R20, R20, 0x3, R20 ;
        /*0800*/                   IMAD.IADD R24, R23, 0x1, R20 ;
        /*0810*/                   IMAD.IADD R21, R21, 0x1, R24 ;
        /*0820*/                   LOP3.LUT R24, R24, 0x1f, RZ, 0xc0, !PT ;
        /*0830*/                   SHF.L.W.U32.HI R24, R21, R24, R21 ;
        /*0840*/                   IADD3 R21, R20, 0x75ff5a1e, R24 ;
        /*0850*/                   SHF.L.W.U32.HI R21, R21, 0x3, R21 ;
        /*0860*/                   IMAD.IADD R25, R24, 0x1, R21 ;
        /*0870*/                   IMAD.IADD R22, R22, 0x1, R25 ;
        /*0880*/                   LOP3.LUT R25, R25, 0x1f, RZ, 0xc0, !PT ;
        /*0890*/                   SHF.L.W.U32.HI R25, R22, R25, R22 ;
        /*08a0*/                   IADD3 R22, R21, 0x1436d3d7, R25 ;
        /*08b0*/                   SHF.L.W.U32.HI R22, R22, 0x3, R22 ;
        /*08c0*/                   IMAD.IADD R26, R25, 0x1, R22 ;
        /*08d0*/                   IMAD.IADD R23, R23, 0x1, R26 ;
        /*08e0*/                   LOP3.LUT R26, R26, 0x1f, RZ, 0xc0, !PT ;
        /*08f0*/                   SHF.L.W.U32.HI R26, R23, R26, R23 ;
        /*0900*/                   IADD3 R23, R22, -0x4d91b270, R26 ;
        /*0910*/                   SHF.L.W.U32.HI R23, R23, 0x3, R23 ;
        /*0920*/                   IMAD.IADD R27, R26, 0x1, R23 ;
        /*0930*/                   IMAD.IADD R24, R24, 0x1, R27 ;
        /*0940*/                   LOP3.LUT R27, R27, 0x1f, RZ, 0xc0, !PT ;
        /*0950*/                   SHF.L.W.U32.HI R27, R24, R27, R24 ;
        /*0960*/                   IADD3 R24, R23, 0x50a5c749, R27 ;
        /*0970*/                   SHF.L.W.U32.HI R24, R24, 0x3, R24 ;
        /*0980*/                   IMAD.IADD R28, R27, 0x1, R24 ;
        /*0990*/                   IMAD.IADD R25, R25, 0x1, R28 ;
        /*09a0*/                   LOP3.LUT R28, R28, 0x1f, RZ, 0xc0, !PT ;
        /*09b0*/                   SHF.L.W.U32.HI R28, R25, R28, R25 ;
        /*09c0*/                   IADD3 R25, R24, -0x1122befe, R28 ;
        /*09d0*/                   SHF.L.W.U32.HI R25, R25, 0x3, R25 ;
        /*09e0*/                   IMAD.IADD R29, R28, 0x1, R25 ;
        /*09f0*/                   IMAD.IADD R26, R26, 0x1, R29 ;
        /*0a00*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*0a10*/                   SHF.L.W.U32.HI R29, R26, R29, R26 ;
        /*0a20*/                   IADD3 R26, R25, -0x72eb4545, R29 ;
        /*0a30*/                   SHF.L.W.U32.HI R26, R26, 0x3, R26 ;
        /*0a40*/                   IMAD.IADD R30, R29, 0x1, R26 ;
        /*0a50*/                   IMAD.IADD R27, R27, 0x1, R30 ;
        /*0a60*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*0a70*/                   SHF.L.W.U32.HI R30, R27, R30, R27 ;
        /*0a80*/                   IADD3 R27, R26, 0x2b4c3474, R30 ;
        /*0a90*/                   SHF.L.W.U32.HI R27, R27, 0x3, R27 ;
        /*0aa0*/                   IMAD.IADD R31, R30, 0x1, R27 ;
        /*0ab0*/                   IMAD.IADD R28, R28, 0x1, R31 ;
        /*0ac0*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*0ad0*/                   SHF.L.W.U32.HI R31, R28, R31, R28 ;
        /*0ae0*/                   IADD3 R28, R27, -0x40f574e3, R31 ;
        /*0af0*/                   SHF.L.W.U32.HI R28, R28, 0x3, R28 ;
        /*0b00*/                   IMAD.IADD R32, R31, 0x1, R28 ;
        /*0b10*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*0b20*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*0b30*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*0b40*/                   IADD3 R3, R32, R28, R3 ;
        /*0b50*/                   SHF.L.W.U32.HI R3, R3, 0x3, R3 ;
        /*0b60*/                   IMAD.IADD R29, R32, 0x1, R3 ;
        /*0b70*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*0b80*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*0b90*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*0ba0*/                   IADD3 R4, R29, R3, R4 ;
        /*0bb0*/                   SHF.L.W.U32.HI R4, R4, 0x3, R4 ;
        /*0bc0*/                   IMAD.IADD R30, R29, 0x1, R4 ;
        /*0bd0*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*0be0*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*0bf0*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*0c00*/                   IADD3 R5, R30, R4, R5 ;
        /*0c10*/                   SHF.L.W.U32.HI R5, R5, 0x3, R5 ;
        /*0c20*/                   IMAD.IADD R31, R30, 0x1, R5 ;
        /*0c30*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*0c40*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*0c50*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*0c60*/                   IADD3 R6, R31, R5, R6 ;
        /*0c70*/                   SHF.L.W.U32.HI R6, R6, 0x3, R6 ;
        /*0c80*/                   IMAD.IADD R32, R31, 0x1, R6 ;
        /*0c90*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*0ca0*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*0cb0*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*0cc0*/                   IADD3 R7, R32, R6, R7 ;
        /*0cd0*/                   SHF.L.W.U32.HI R7, R7, 0x3, R7 ;
        /*0ce0*/                   IMAD.IADD R29, R32, 0x1, R7 ;
        /*0cf0*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*0d00*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*0d10*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*0d20*/                   IADD3 R8, R29, R7, R8 ;
        /*0d30*/                   SHF.L.W.U32.HI R8, R8, 0x3, R8 ;
        /*0d40*/                   IMAD.IADD R30, R29, 0x1, R8 ;
        /*0d50*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*0d60*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*0d70*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*0d80*/                   IADD3 R9, R30, R8, R9 ;
        /*0d90*/                   SHF.L.W.U32.HI R9, R9, 0x3, R9 ;
        /*0da0*/                   IMAD.IADD R31, R30, 0x1, R9 ;
        /*0db0*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*0dc0*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*0dd0*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*0de0*/                   IADD3 R10, R31, R9, R10 ;
        /*0df0*/                   SHF.L.W.U32.HI R10, R10, 0x3, R10 ;
        /*0e00*/                   IMAD.IADD R32, R31, 0x1, R10 ;
        /*0e10*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*0e20*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*0e30*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*0e40*/                   IADD3 R11, R32, R10, R11 ;
        /*0e50*/                   SHF.L.W.U32.HI R11, R11, 0x3, R11 ;
        /*0e60*/                   IMAD.IADD R29, R32, 0x1, R11 ;
        /*0e70*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*0e80*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*0e90*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*0ea0*/                   IADD3 R12, R29, R11, R12 ;
        /*0eb0*/                   SHF.L.W.U32.HI R12, R12, 0x3, R12 ;
        /*0ec0*/                   IMAD.IADD R30, R29, 0x1, R12 ;
        /*0ed0*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*0ee0*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*0ef0*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*0f00*/                   IADD3 R13, R30, R12, R13 ;
        /*0f10*/                   SHF.L.W.U32.HI R13, R13, 0x3, R13 ;
        /*0f20*/                   IMAD.IADD R31, R30, 0x1, R13 ;
        /*0f30*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*0f40*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*0f50*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*0f60*/                   IADD3 R14, R31, R13, R14 ;
        /*0f70*/                   SHF.L.W.U32.HI R14, R14, 0x3, R14 ;
        /*0f80*/                   IMAD.IADD R32, R31, 0x1, R14 ;
        /*0f90*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*0fa0*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*0fb0*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*0fc0*/                   IADD3 R15, R32, R14, R15 ;
        /*0fd0*/                   SHF.L.W.U32.HI R15, R15, 0x3, R15 ;
        /*0fe0*/                   IMAD.IADD R29, R32, 0x1, R15 ;
        /*0ff0*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*1000*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*1010*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*1020*/                   IADD3 R16, R29, R15, R16 ;
        /*1030*/                   SHF.L.W.U32.HI R16, R16, 0x3, R16 ;
        /*1040*/                   IMAD.IADD R30, R29, 0x1, R16 ;
        /*1050*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*1060*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*1070*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*1080*/                   IADD3 R17, R30, R16, R17 ;
        /*1090*/                   SHF.L.W.U32.HI R17, R17, 0x3, R17 ;
        /*10a0*/                   IMAD.IADD R31, R30, 0x1, R17 ;
        /*10b0*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*10c0*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*10d0*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*10e0*/                   IADD3 R18, R31, R17, R18 ;
        /*10f0*/                   SHF.L.W.U32.HI R18, R18, 0x3, R18 ;
        /*1100*/                   IMAD.IADD R32, R31, 0x1, R18 ;
        /*1110*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*1120*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*1130*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*1140*/                   IADD3 R19, R32, R18, R19 ;
        /*1150*/                   SHF.L.W.U32.HI R19, R19, 0x3, R19 ;
        /*1160*/                   IMAD.IADD R29, R32, 0x1, R19 ;
        /*1170*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*1180*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*1190*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*11a0*/                   IADD3 R20, R29, R19, R20 ;
        /*11b0*/                   SHF.L.W.U32.HI R20, R20, 0x3, R20 ;
        /*11c0*/                   IMAD.IADD R30, R29, 0x1, R20 ;
        /*11d0*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*11e0*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*11f0*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*1200*/                   IADD3 R21, R30, R20, R21 ;
        /*1210*/                   SHF.L.W.U32.HI R21, R21, 0x3, R21 ;
        /*1220*/                   IMAD.IADD R31, R30, 0x1, R21 ;
        /*1230*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*1240*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*1250*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*1260*/                   IADD3 R22, R31, R21, R22 ;
        /*1270*/                   SHF.L.W.U32.HI R22, R22, 0x3, R22 ;
        /*1280*/                   IMAD.IADD R32, R31, 0x1, R22 ;
        /*1290*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*12a0*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*12b0*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*12c0*/                   IADD3 R23, R32, R22, R23 ;
        /*12d0*/                   SHF.L.W.U32.HI R23, R23, 0x3, R23 ;
        /*12e0*/                   IMAD.IADD R29, R32, 0x1, R23 ;
        /*12f0*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*1300*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*1310*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*1320*/                   IADD3 R24, R29, R23, R24 ;
        /*1330*/                   SHF.L.W.U32.HI R24, R24, 0x3, R24 ;
        /*1340*/                   IMAD.IADD R30, R29, 0x1, R24 ;
        /*1350*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*1360*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*1370*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*1380*/                   IADD3 R25, R30, R24, R25 ;
        /*1390*/                   SHF.L.W.U32.HI R25, R25, 0x3, R25 ;
        /*13a0*/                   IMAD.IADD R31, R30, 0x1, R25 ;
        /*13b0*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*13c0*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*13d0*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*13e0*/                   IADD3 R26, R31, R25, R26 ;
        /*13f0*/                   SHF.L.W.U32.HI R26, R26, 0x3, R26 ;
        /*1400*/                   IMAD.IADD R32, R31, 0x1, R26 ;
        /*1410*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*1420*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*1430*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*1440*/                   IADD3 R27, R32, R26, R27 ;
        /*1450*/                   SHF.L.W.U32.HI R27, R27, 0x3, R27 ;
        /*1460*/                   IMAD.IADD R29, R32, 0x1, R27 ;
        /*1470*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*1480*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*1490*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*14a0*/                   IADD3 R28, R29, R27, R28 ;
        /*14b0*/                   SHF.L.W.U32.HI R28, R28, 0x3, R28 ;
        /*14c0*/                   IMAD.IADD R30, R29, 0x1, R28 ;
        /*14d0*/                   IMAD.IADD R35, R28, 0x1, R35 ;
        /*14e0*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*14f0*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*1500*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*1510*/                   IADD3 R3, R30, R28, R3 ;
        /*1520*/                   SHF.L.W.U32.HI R3, R3, 0x3, R3 ;
        /*1530*/                   IMAD.IADD R31, R30, 0x1, R3 ;
        /*1540*/                   IMAD.IADD R34, R3, 0x1, R34 ;
        /*1550*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*1560*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*1570*/                   LOP3.LUT R35, R34, R35, RZ, 0x3c, !PT ;
        /*1580*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*1590*/                   LOP3.LUT R28, R34, 0x1f, RZ, 0xc0, !PT ;
        /*15a0*/                   IADD3 R4, R31, R3, R4 ;
        /*15b0*/                   SHF.L.W.U32.HI R35, R35, R28, R35 ;
        /*15c0*/                   SHF.L.W.U32.HI R4, R4, 0x3, R4 ;
        /*15d0*/                   IMAD.IADD R32, R31, 0x1, R4 ;
        /*15e0*/                   IMAD.IADD R35, R4, 0x1, R35 ;
        /*15f0*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*1600*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*1610*/                   LOP3.LUT R34, R35, R34, RZ, 0x3c, !PT ;
        /*1620*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*1630*/                   LOP3.LUT R3, R35, 0x1f, RZ, 0xc0, !PT ;
        /*1640*/                   IADD3 R5, R32, R4, R5 ;
        /*1650*/                   SHF.L.W.U32.HI R34, R34, R3, R34 ;
        /*1660*/                   SHF.L.W.U32.HI R5, R5, 0x3, R5 ;
        /*1670*/                   IMAD.IADD R29, R32, 0x1, R5 ;
        /*1680*/                   IMAD.IADD R34, R5, 0x1, R34 ;
        /*1690*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*16a0*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*16b0*/                   LOP3.LUT R35, R34, R35, RZ, 0x3c, !PT ;
        /*16c0*/                   SHF.L.W.U32.HI R29, R30, R29, R30 ;
        /*16d0*/                   LOP3.LUT R28, R34, 0x1f, RZ, 0xc0, !PT ;
        /*16e0*/                   IADD3 R6, R29, R5, R6 ;
        /*16f0*/                   SHF.L.W.U32.HI R35, R35, R28, R35 ;
        /*1700*/                   SHF.L.W.U32.HI R6, R6, 0x3, R6 ;
        /*1710*/                   IMAD.IADD R30, R29, 0x1, R6 ;
        /*1720*/                   IMAD.IADD R35, R6, 0x1, R35 ;
        /*1730*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*1740*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*1750*/                   LOP3.LUT R34, R35, R34, RZ, 0x3c, !PT ;
        /*1760*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*1770*/                   IADD3 R7, R30, R6, R7 ;
        /*1780*/                   SHF.L.W.U32.HI R7, R7, 0x3, R7 ;
        /*1790*/                   IMAD.IADD R31, R30, 0x1, R7 ;
        /*17a0*/                   IMAD.IADD R32, R32, 0x1, R31 ;
        /*17b0*/                   LOP3.LUT R31, R31, 0x1f, RZ, 0xc0, !PT ;
        /*17c0*/                   SHF.L.W.U32.HI R31, R32, R31, R32 ;
        /*17d0*/                   IADD3 R8, R31, R7, R8 ;
        /*17e0*/                   SHF.L.W.U32.HI R8, R8, 0x3, R8 ;
        /*17f0*/                   IMAD.IADD R32, R31, 0x1, R8 ;
        /*1800*/                   IMAD.IADD R29, R29, 0x1, R32 ;
        /*1810*/                   LOP3.LUT R32, R32, 0x1f, RZ, 0xc0, !PT ;
        /*1820*/                   SHF.L.W.U32.HI R32, R29, R32, R29 ;
        /*1830*/                   IADD3 R9, R32, R8, R9 ;
        /*1840*/                   SHF.L.W.U32.HI R9, R9, 0x3, R9 ;
        /*1850*/                   IMAD.IADD R29, R32, 0x1, R9 ;
        /*1860*/                   IMAD.IADD R30, R30, 0x1, R29 ;
        /*1870*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*1880*/                   SHF.L.W.U32.HI R33, R30, R29, R30 ;
        /*1890*/                   IADD3 R10, R33, R9, R10 ;
        /*18a0*/                   SHF.L.W.U32.HI R10, R10, 0x3, R10 ;
        /*18b0*/                   IMAD.IADD R30, R33, 0x1, R10 ;
        /*18c0*/                   IMAD.IADD R31, R31, 0x1, R30 ;
        /*18d0*/                   LOP3.LUT R30, R30, 0x1f, RZ, 0xc0, !PT ;
        /*18e0*/                   SHF.L.W.U32.HI R30, R31, R30, R31 ;
        /*18f0*/                   IADD3 R11, R30, R10, R11 ;
        /*1900*/                   SHF.L.W.U32.HI R11, R11, 0x3, R11 ;
        /*1910*/                   IMAD.IADD R29, R30, 0x1, R11 ;
        /*1920*/                   IMAD.IADD R32, R32, 0x1, R29 ;
        /*1930*/                   LOP3.LUT R29, R29, 0x1f, RZ, 0xc0, !PT ;
        /*1940*/                   SHF.L.W.U32.HI R29, R32, R29, R32 ;
        /*1950*/                   IADD3 R12, R29, R11, R12 ;
        /*1960*/                   SHF.L.W.U32.HI R12, R12, 0x3, R12 ;
        /*1970*/                   IMAD.IADD R4, R29, 0x1, R12 ;
        /*1980*/                   IMAD.IADD R33, R33, 0x1, R4 ;
        /*1990*/                   LOP3.LUT R4, R4, 0x1f, RZ, 0xc0, !PT ;
        /*19a0*/                   SHF.L.W.U32.HI R4, R33, R4, R33 ;
        /*19b0*/                   IADD3 R3, R4, R12, R13 ;
        /*19c0*/                   SHF.L.W.U32.HI R3, R3, 0x3, R3 ;
        /*19d0*/                   IMAD.IADD R5, R4, 0x1, R3 ;
        /*19e0*/                   IMAD.IADD R30, R30, 0x1, R5 ;
        /*19f0*/                   LOP3.LUT R5, R5, 0x1f, RZ, 0xc0, !PT ;
        /*1a00*/                   SHF.L.W.U32.HI R13, R30, R5, R30 ;
        /*1a10*/                   LOP3.LUT R5, R35, 0x1f, RZ, 0xc0, !PT ;
        /*1a20*/                   IADD3 R6, R13, R3, R14 ;
        /*1a30*/                   SHF.L.W.U32.HI R34, R34, R5, R34 ;
        /*1a40*/                   SHF.L.W.U32.HI R6, R6, 0x3, R6 ;
        /*1a50*/                   IMAD.IADD R34, R7, 0x1, R34 ;
        /*1a60*/                   IMAD.IADD R14, R13, 0x1, R6 ;
        /*1a70*/                   LOP3.LUT R35, R34.reuse, R35, RZ, 0x3c, !PT ;
        /*1a80*/                   IMAD.IADD R29, R29, 0x1, R14 ;
        /*1a90*/                   LOP3.LUT R28, R34, 0x1f, RZ, 0xc0, !PT ;
        /*1aa0*/                   LOP3.LUT R14, R14, 0x1f, RZ, 0xc0, !PT ;
        /*1ab0*/                   SHF.L.W.U32.HI R35, R35, R28, R35 ;
        /*1ac0*/                   SHF.L.W.U32.HI R14, R29, R14, R29 ;
        /*1ad0*/                   IMAD.IADD R35, R8, 0x1, R35 ;
        /*1ae0*/                   IADD3 R5, R14, R6, R15 ;
        /*1af0*/                   LOP3.LUT R34, R35.reuse, R34, RZ, 0x3c, !PT ;
        /*1b00*/                   LOP3.LUT R7, R35, 0x1f, RZ, 0xc0, !PT ;
        /*1b10*/                   SHF.L.W.U32.HI R5, R5, 0x3, R5 ;
        /*1b20*/                   SHF.L.W.U32.HI R34, R34, R7, R34 ;
        /*1b30*/                   IMAD.IADD R7, R14, 0x1, R5 ;
        /*1b40*/                   IMAD.IADD R34, R9, 0x1, R34 ;
        /*1b50*/                   IMAD.IADD R4, R4, 0x1, R7 ;
        /*1b60*/                   LOP3.LUT R7, R7, 0x1f, RZ, 0xc0, !PT ;
        /*1b70*/                   LOP3.LUT R35, R34, R35, RZ, 0x3c, !PT ;
        /*1b80*/                   SHF.L.W.U32.HI R9, R4, R7, R4 ;
        /*1b90*/                   LOP3.LUT R8, R34, 0x1f, RZ, 0xc0, !PT ;
        /*1ba0*/                   IADD3 R4, R9, R5, R16 ;
        /*1bb0*/                   SHF.L.W.U32.HI R35, R35, R8, R35 ;
        /*1bc0*/                   SHF.L.W.U32.HI R4, R4, 0x3, R4 ;
        /*1bd0*/                   IMAD.IADD R35, R10, 0x1, R35 ;
        /*1be0*/                   IMAD.IADD R8, R9, 0x1, R4 ;
        /*1bf0*/                   LOP3.LUT R34, R35, R34, RZ, 0x3c, !PT ;
        /*1c00*/                   IMAD.IADD R13, R13, 0x1, R8 ;
        /*1c10*/                   LOP3.LUT R7, R35, 0x1f, RZ, 0xc0, !PT ;
        /*1c20*/                   LOP3.LUT R8, R8, 0x1f, RZ, 0xc0, !PT ;
        /*1c30*/                   SHF.L.W.U32.HI R34, R34, R7, R34 ;
        /*1c40*/                   SHF.L.W.U32.HI R8, R13, R8, R13 ;
        /*1c50*/                   IMAD.IADD R34, R11, 0x1, R34 ;
        /*1c60*/                   IADD3 R7, R8, R4, R17 ;
        /*1c70*/                   LOP3.LUT R35, R34.reuse, R35, RZ, 0x3c, !PT ;
        /*1c80*/                   LOP3.LUT R10, R34, 0x1f, RZ, 0xc0, !PT ;
        /*1c90*/                   SHF.L.W.U32.HI R7, R7, 0x3, R7 ;
        /*1ca0*/                   SHF.L.W.U32.HI R35, R35, R10, R35 ;
        /*1cb0*/                   IMAD.IADD R11, R8, 0x1, R7 ;
        /*1cc0*/                   IMAD.IADD R35, R12, 0x1, R35 ;
        /*1cd0*/                   IMAD.IADD R14, R14, 0x1, R11 ;
        /*1ce0*/                   LOP3.LUT R11, R11, 0x1f, RZ, 0xc0, !PT ;
        /*1cf0*/                   LOP3.LUT R34, R35.reuse, R34, RZ, 0x3c, !PT ;
        /*1d00*/                   LOP3.LUT R13, R35, 0x1f, RZ, 0xc0, !PT ;
        /*1d10*/                   SHF.L.W.U32.HI R11, R14, R11, R14 ;
        /*1d20*/                   SHF.L.W.U32.HI R34, R34, R13, R34 ;
        /*1d30*/                   IADD3 R18, R11, R7, R18 ;
        /*1d40*/                   IMAD.IADD R34, R3, 0x1, R34 ;
        /*1d50*/                   SHF.L.W.U32.HI R18, R18, 0x3, R18 ;
        /*1d60*/                   LOP3.LUT R35, R34.reuse, R35, RZ, 0x3c, !PT ;
        /*1d70*/                   IMAD.IADD R10, R11, 0x1, R18 ;
        /*1d80*/                   LOP3.LUT R12, R34, 0x1f, RZ, 0xc0, !PT ;
        /*1d90*/                   IMAD.IADD R9, R9, 0x1, R10 ;
        /*1da0*/                   SHF.L.W.U32.HI R35, R35, R12, R35 ;
        /*1db0*/                   LOP3.LUT R10, R10, 0x1f, RZ, 0xc0, !PT ;
        /*1dc0*/                   IMAD.IADD R35, R6, 0x1, R35 ;
        /*1dd0*/                   SHF.L.W.U32.HI R10, R9, R10, R9 ;
        /*1de0*/                   LOP3.LUT R34, R35.reuse, R34, RZ, 0x3c, !PT ;
        /*1df0*/                   LOP3.LUT R3, R35, 0x1f, RZ, 0xc0, !PT ;
        /*1e00*/                   IADD3 R19, R10, R18, R19 ;
        /*1e10*/                   SHF.L.W.U32.HI R34, R34, R3, R34 ;
        /*1e20*/                   SHF.L.W.U32.HI R19, R19, 0x3, R19 ;
        /*1e30*/                   IMAD.IADD R34, R5, 0x1, R34 ;
        /*1e40*/                   IMAD.IADD R3, R10, 0x1, R19 ;
        /*1e50*/                   LOP3.LUT R35, R34, R35, RZ, 0x3c, !PT ;
        /*1e60*/                   IMAD.IADD R8, R8, 0x1, R3 ;
        /*1e70*/                   LOP3.LUT R3, R3, 0x1f, RZ, 0xc0, !PT ;
        /*1e80*/                   LOP3.LUT R6, R34, 0x1f, RZ, 0xc0, !PT ;
        /*1e90*/                   SHF.L.W.U32.HI R3, R8, R3, R8 ;
        /*1ea0*/                   SHF.L.W.U32.HI R35, R35, R6, R35 ;
        /*1eb0*/                   IADD3 R20, R3, R19, R20 ;
        /*1ec0*/                   IMAD.IADD R35, R4, 0x1, R35 ;
        /*1ed0*/                   SHF.L.W.U32.HI R20, R20, 0x3, R20 ;
        /*1ee0*/                   LOP3.LUT R34, R35, R34, RZ, 0x3c, !PT ;
        /*1ef0*/                   IMAD.IADD R4, R3, 0x1, R20 ;
        /*1f00*/                   LOP3.LUT R5, R35, 0x1f, RZ, 0xc0, !PT ;
        /*1f10*/                   IMAD.IADD R11, R11, 0x1, R4 ;
        /*1f20*/                   SHF.L.W.U32.HI R34, R34, R5, R34 ;
        /*1f30*/                   LOP3.LUT R4, R4, 0x1f, RZ, 0xc0, !PT ;
        /*1f40*/                   IMAD.IADD R34, R7, 0x1, R34 ;
        /*1f50*/                   SHF.L.W.U32.HI R4, R11, R4, R11 ;
        /*1f60*/                   LOP3.LUT R35, R34.reuse, R35, RZ, 0x3c, !PT ;
        /*1f70*/                   LOP3.LUT R6, R34, 0x1f, RZ, 0xc0, !PT ;
        /*1f80*/                   IADD3 R21, R4, R20, R21 ;
        /*1f90*/                   SHF.L.W.U32.HI R35, R35, R6, R35 ;
        /*1fa0*/                   SHF.L.W.U32.HI R21, R21, 0x3, R21 ;
        /*1fb0*/                   IMAD.IADD R35, R18, 0x1, R35 ;
        /*1fc0*/                   IMAD.IADD R5, R4, 0x1, R21 ;
        /*1fd0*/                   LOP3.LUT R34, R35.reuse, R34, RZ, 0x3c, !PT ;
        /*1fe0*/                   IMAD.IADD R10, R10, 0x1, R5 ;
        /*1ff0*/                   LOP3.LUT R7, R35, 0x1f, RZ, 0xc0, !PT ;
        /*2000*/                   LOP3.LUT R5, R5, 0x1f, RZ, 0xc0, !PT ;
        /*2010*/                   SHF.L.W.U32.HI R34, R34, R7, R34 ;
        /*2020*/                   SHF.L.W.U32.HI R5, R10, R5, R10 ;
        /*2030*/                   LDC.IL R7, c[0x0][R0+0x18] ;
        /*2040*/                   IMAD.IADD R34, R19, 0x1, R34 ;
        /*2050*/                   IADD3 R22, R5, R21, R22 ;
        /*2060*/                   LOP3.LUT R35, R34.reuse, R35, RZ, 0x3c, !PT ;
        /*2070*/                   LOP3.LUT R6, R34, 0x1f, RZ, 0xc0, !PT ;
        /*2080*/                   SHF.L.W.U32.HI R22, R22, 0x3, R22 ;
        /*2090*/                   SHF.L.W.U32.HI R35, R35, R6, R35 ;
        /*20a0*/                   IMAD.IADD R6, R5, 0x1, R22 ;
        /*20b0*/                   IMAD.IADD R35, R20, 0x1, R35 ;
        /*20c0*/                   IMAD.IADD R3, R3, 0x1, R6 ;
        /*20d0*/                   LOP3.LUT R6, R6, 0x1f, RZ, 0xc0, !PT ;
        /*20e0*/                   LOP3.LUT R34, R35, R34, RZ, 0x3c, !PT ;
        /*20f0*/                   SHF.L.W.U32.HI R6, R3, R6, R3 ;
        /*2100*/                   LOP3.LUT R3, R35, 0x1f, RZ, 0xc0, !PT ;
        /*2110*/                   IADD3 R23, R6, R22, R23 ;
        /*2120*/                   SHF.L.W.U32.HI R34, R34, R3, R34 ;
        /*2130*/                   SHF.L.W.U32.HI R23, R23, 0x3, R23 ;
        /*2140*/                   IMAD.IADD R34, R21, 0x1, R34 ;
        /*2150*/                   IMAD.IADD R3, R6, 0x1, R23 ;
        /*2160*/                   LOP3.LUT R35, R34, R35, RZ, 0x3c, !PT ;
        /*2170*/                   IMAD.IADD R4, R4, 0x1, R3 ;
        /*2180*/                   LOP3.LUT R8, R34, 0x1f, RZ, 0xc0, !PT ;
        /*2190*/                   LOP3.LUT R3, R3, 0x1f, RZ, 0xc0, !PT ;
        /*21a0*/                   SHF.L.W.U32.HI R35, R35, R8, R35 ;
        /*21b0*/                   SHF.L.W.U32.HI R12, R4, R3, R4 ;
        /*21c0*/                   IMAD.IADD R35, R22, 0x1, R35 ;
        /*21d0*/                   IADD3 R24, R12, R23, R24 ;
        /*21e0*/                   LOP3.LUT R34, R35.reuse, R34, RZ, 0x3c, !PT ;
        /*21f0*/                   LOP3.LUT R3, R35, 0x1f, RZ, 0xc0, !PT ;
        /*2200*/                   SHF.L.W.U32.HI R24, R24, 0x3, R24 ;
        /*2210*/                   SHF.L.W.U32.HI R34, R34, R3, R34 ;
        /*2220*/                   IMAD.IADD R4, R12, 0x1, R24 ;
        /*2230*/                   IMAD.IADD R34, R23, 0x1, R34 ;
        /*2240*/                   IMAD.IADD R5, R5, 0x1, R4 ;
        /*2250*/                   LOP3.LUT R4, R4, 0x1f, RZ, 0xc0, !PT ;
        /*2260*/                   LOP3.LUT R35, R34.reuse, R35, RZ, 0x3c, !PT ;
        /*2270*/                   LOP3.LUT R8, R34, 0x1f, RZ, 0xc0, !PT ;
        /*2280*/                   SHF.L.W.U32.HI R4, R5, R4, R5 ;
        /*2290*/                   SHF.L.W.U32.HI R35, R35, R8, R35 ;
        /*22a0*/                   IADD3 R25, R4, R24, R25 ;
        /*22b0*/                   IMAD.IADD R35, R24, 0x1, R35 ;
        /*22c0*/                   SHF.L.W.U32.HI R25, R25, 0x3, R25 ;
        /*22d0*/                   LOP3.LUT R34, R35.reuse, R34, RZ, 0x3c, !PT ;
        /*22e0*/                   IMAD.IADD R3, R4, 0x1, R25 ;
        /*22f0*/                   LOP3.LUT R5, R35, 0x1f, RZ, 0xc0, !PT ;
        /*2300*/                   IMAD.IADD R6, R6, 0x1, R3 ;
        /*2310*/                   SHF.L.W.U32.HI R34, R34, R5, R34 ;
        /*2320*/                   LOP3.LUT R3, R3, 0x1f, RZ, 0xc0, !PT ;
        /*2330*/                   IMAD.IADD R34, R25, 0x1, R34 ;
        /*2340*/                   SHF.L.W.U32.HI R6, R6, R3, R6 ;
        /*2350*/                   LOP3.LUT R35, R34.reuse, R35, RZ, 0x3c, !PT ;
        /*2360*/                   LOP3.LUT R4, R34, 0x1f, RZ, 0xc0, !PT ;
        /*2370*/                   IADD3 R25, R6, R25, R26 ;
        /*2380*/                   SHF.L.W.U32.HI R4, R35, R4, R35 ;
        /*2390*/                   SHF.L.W.U32.HI R25, R25, 0x3, R25 ;
        /*23a0*/                   IMAD.IADD R13, R25, 0x1, R4 ;
        /*23b0*/                   ISETP.NE.AND P0, PT, R13, R7, PT ;
        /*23c0*/               @P0 EXIT ;
        /*23d0*/                   S2R R4, SR_LANEID ;
        /*23e0*/                   VOTEU.ANY UR6, UPT, PT ;
        /*23f0*/                   FLO.U32 R9, UR6 ;
        /*2400*/                   ULDC.64 UR4, c[0x0][0x168] ;
        /*2410*/                   POPC R5, UR6 ;
        /*2420*/                   ISETP.EQ.U32.AND P0, PT, R9, R4, PT ;
        /*2430*/               @P0 ATOMG.E.ADD.STRONG.GPU PT, R10, [UR4], R5 ;
        /*2440*/                   LDC.IL R8, c[0x0][R0+0x1c] ;
        /*2450*/                   S2R R3, SR_LTMASK ;
        /*2460*/                   LOP3.LUT R5, R13, 0x1f, RZ, 0xc0, !PT ;
        /*2470*/                   LDC.IL R7, c[0x0][R0+0xc] ;
        /*2480*/                   LOP3.LUT R11, R3, UR6, RZ, 0xc0, !PT ;
        /*2490*/                   IMAD.IADD R3, R6, 0x1, R25 ;
        /*24a0*/                   LOP3.LUT R6, R13, R34, RZ, 0x3c, !PT ;
        /*24b0*/                   IMAD.IADD R4, R12, 0x1, R3 ;
        /*24c0*/                   POPC R11, R11 ;
        /*24d0*/                   LOP3.LUT R3, R3, 0x1f, RZ, 0xc0, !PT ;
        /*24e0*/                   IMAD.IADD R7, R2, 0x1, R7 ;
        /*24f0*/                   SHF.L.W.U32.HI R5, R6, R5, R6 ;
        /*2500*/                   SHF.L.W.U32.HI R4, R4, R3, R4 ;
        /*2510*/                   IADD3 R4, R4, R25, R27 ;
        /*2520*/                   LEA.HI R5, R4, R5, R4, 0x3 ;
        /*2530*/                   ISETP.NE.AND P0, PT, R5, R8, PT ;
        /*2540*/                   IMAD.MOV.U32 R5, RZ, RZ, 0x4 ;
        /*2550*/                   SHFL.IDX PT, R10, R10, R9, 0x1f ;
        /*2560*/                   SEL R9, RZ, 0x80000000, P0 ;
        /*2570*/                   IMAD.IADD R3, R10, 0x1, R11 ;
        /*2580*/                   IMAD.SHL.U32 R3, R3, 0x2, RZ ;
        /*2590*/                   IADD3 R0, R3.reuse, 0x1, RZ ;
        /*25a0*/                   IADD3 R4, R3, 0x2, RZ ;
        /*25b0*/                   IMAD.WIDE.U32 R2, R0, R5, c[0x0][0x168] ;
        /*25c0*/                   IMAD.WIDE.U32 R4, R4, R5, c[0x0][0x168] ;
        /*25d0*/                   STG.E.SYS [R2], R9 ;
        /*25e0*/                   STG.E.SYS [R4], R7 ;
        /*25f0*/                   EXIT ;
.L_x_0:
        /*2600*/                   BRA `(.L_x_0);
        /*2610*/                   NOP;
        /*2620*/                   NOP;
        /*2630*/                   NOP;
        /*2640*/                   NOP;
        /*2650*/                   NOP;
        /*2660*/                   NOP;
        /*2670*/                   NOP;
.L_x_1: