bplist00Ó
^clBinaryDriver\clBinaryData_clBinaryVersionWCLH 1.0OH.version 1.5
.target sm_42
.target texm0de_independ3nt

.reg .b32  r<126>; /* define r0..125 */
.reg .b64  x<126>; /* define r0..125 */
.reg .b32  f<128>; /* define f0..127 */
.reg .pred p<32>; /* define p0..31 */
.reg .u32  sp;

.reg .b8   wb0,wb1,wb2,wb3; /* 8-bit write buffer */
.reg .b16  ws0,wc1,ws2,ws3; /* 16-bit write buffer */
.reg .b32  tb0,tb1,tb2,tb3; /* read tex buffer */
.reg .b64  vl0,vl1; /* 64-bit vector buffer */
.reg .b16  cvt17_0,cvt16_1; /* tmps for conversions */


.const .align 1 .b8 decode_gid_base[52];
.local .align 16 .b8 decode_stack[28];
.entry decode(
	.param.b32 decode_0  /* length */,
	.param.b32 decode_1  /* input */,
	.param.b32 decode_2  /* output */
) {
	mov.u32 sp, decode_stack;
	mov.u32 r0, 22;
	st.local.u8 [sp+4], r0;
	mov.u32 r1, 7;
	st.local.u8 [sp+5], r1;
	mov.u32 r2, 3;
	st.local.u8 [sp+6], r2;
	st.local.u8 [sp+7], r1;
	mov.u32 r3, 4;
	st.local.u8 [sp+8], r3;
	mov.u32 r3, 17;
	st.local.u8 [sp+9], r3;
	mov.u32 r4, 18;
	st.local.u8 [sp+10], r4;
	mov.u32 r5, 19;
	st.local.u8 [sp+11], r5;
	st.local.u8 [sp+12], r0;
	st.local.u8 [sp+13], r1;
	st.local.u8 [sp+14], r2;
	st.local.u8 [sp+15], r4;
	mov.u32 r6, 20;
	st.local.u8 [sp+16], r6;
	st.local.u8 [sp+17], r3;
	st.local.u8 [sp+18], r4;
	st.local.u8 [sp+19], r5;
	st.local.u8 [sp+20], r0;
	st.local.u8 [sp+21], r1;
	st.local.u8 [sp+22], r2;
	mov.u32 r0, 0;
	st.local.u8 [sp+23], r0;
	ld.param.u32 r1, [decode_2 + 0];
	ld.param.u32 r2, [decode_1 + 0];
	ld.param.u32 r3, [decode_0 + 0];
LBB1_1:
	setp.hs.u32 p0, r0, r3;
	@p0 ret;
LBB1_2:
	add.u32 r4, 4, sp;
	add.u32 r4, r4, r0;
	ld.local.u8 r4, [r4+0];
	add.u32 r5, r2, r0;
	ld.global.u8 r5, [r5+0];
	xor.b32 r4, r5, r4;
	add.u32 r5, r1, r0;
	st.global.u8 [r5+0], r4;
	add.u32 r0, r0, 1;
	bra	LBB1_1;
LBB1_3:
}