Add files using upload-large-folder tool

.gitignore

@@ -0,0 +1,2 @@
+counter.pkl
+*.cu

config.json

@@ -0,0 +1,3 @@
+{
+  "dfloat11_version": "0.1.0"
+}

counter_original.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:b6a8b405b0f5d58b264a6902d9975af85df4c6002341a04aa23a73aafaee39ac
+size 270

decode.ptx

@@ -0,0 +1,539 @@
+//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: CL-34097967
+// Cuda compilation tools, release 12.4, V12.4.131
+// Based on NVVM 7.0.1
+//
+
+.version 8.4
+.target sm_52
+.address_size 64
+
+	// .globl	huffman_decode
+.const .align 1 .b8 shared_mem_constants[1280] = {120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 120, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 116, 115, 115, 115, 115, 115, 115, 115, 115, 115, 115, 115, 115, 115, 115, 115, 115, 114, 114, 114, 114, 114, 114, 114, 114, 113, 113, 113, 113, 112, 112, 111, 0, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 117, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 118, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 119, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 109, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 108, 107, 107, 107, 107, 107, 107, 107, 107, 107, 107, 107, 107, 107, 107, 107, 107, 106, 106, 106, 106, 106, 106, 106, 106, 104, 104, 103, 0, 122, 122, 105, 105, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 121, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 110, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 102, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 101, 100, 100, 100, 100, 100, 100, 100, 100, 100, 100, 100, 100, 100, 100, 100, 100, 99, 99, 99, 99, 99, 99, 99, 99, 98, 98, 98, 98, 97, 97, 124, 0, 91, 91, 92, 92, 125, 125, 126, 126, 0, 81, 84, 85, 87, 88, 89, 90, 93, 93, 93, 93, 93, 93, 93, 93, 93, 93, 93, 93, 93, 93, 93, 93, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 94, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 95, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 96, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 32, 0, 0, 32, 32, 0, 32, 32, 32, 32, 31, 31, 28, 27, 26, 25, 23, 22, 21, 20, 19, 18, 16, 15, 15, 13, 12, 11, 10, 10, 8, 7, 6, 5, 4, 3, 3, 2, 2, 3, 10, 15, 17, 24, 31, 31};
+.extern .shared .align 16 .b8 shared_mem[];
+
+.visible .entry huffman_decode(
+	.param .u64 huffman_decode_param_0,
+	.param .u64 huffman_decode_param_1,
+	.param .u64 huffman_decode_param_2,
+	.param .u64 huffman_decode_param_3,
+	.param .u64 huffman_decode_param_4,
+	.param .u32 huffman_decode_param_5,
+	.param .u32 huffman_decode_param_6,
+	.param .u32 huffman_decode_param_7
+)
+{
+	.reg .pred 	%p<33>;
+	.reg .b16 	%rs<48>;
+	.reg .b32 	%r<253>;
+	.reg .b64 	%rd<102>;
+
+
+	ld.param.u64 	%rd26, [huffman_decode_param_0];
+	ld.param.u64 	%rd22, [huffman_decode_param_1];
+	ld.param.u64 	%rd23, [huffman_decode_param_2];
+	ld.param.u64 	%rd24, [huffman_decode_param_3];
+	ld.param.u64 	%rd25, [huffman_decode_param_4];
+	ld.param.u32 	%r65, [huffman_decode_param_5];
+	ld.param.u32 	%r66, [huffman_decode_param_6];
+	ld.param.u32 	%r67, [huffman_decode_param_7];
+	cvta.to.global.u64 	%rd1, %rd26;
+	mov.u32 	%r1, %ntid.x;
+	mul.lo.s32 	%r2, %r1, %r65;
+	mov.u32 	%r68, shared_mem;
+	add.s32 	%r3, %r68, %r2;
+	shl.b32 	%r69, %r1, 2;
+	add.s32 	%r4, %r2, %r69;
+	mov.u32 	%r5, %ctaid.x;
+	mov.u32 	%r252, %tid.x;
+	mad.lo.s32 	%r7, %r5, %r1, %r252;
+	mul.lo.s32 	%r8, %r252, %r65;
+	setp.gt.s32 	%p1, %r252, 1279;
+	@%p1 bra 	$L__BB0_3;
+
+	mov.u32 	%r235, %r252;
+
+$L__BB0_2:
+	cvt.s64.s32 	%rd27, %r235;
+	mov.u64 	%rd28, shared_mem_constants;
+	add.s64 	%rd29, %rd28, %rd27;
+	ld.const.u8 	%rs19, [%rd29];
+	add.s32 	%r71, %r68, %r235;
+	st.volatile.shared.u8 	[%r71], %rs19;
+	add.s32 	%r235, %r235, %r1;
+	setp.lt.s32 	%p2, %r235, 1280;
+	@%p2 bra 	$L__BB0_2;
+
+$L__BB0_3:
+	add.s32 	%r11, %r1, -1;
+	setp.eq.s32 	%p3, %r252, %r11;
+	add.s32 	%r72, %r65, 4;
+	selp.b32 	%r12, %r72, %r65, %p3;
+	mul.lo.s32 	%r13, %r7, %r65;
+	sub.s32 	%r73, %r66, %r13;
+	min.s32 	%r14, %r12, %r73;
+	setp.gt.s32 	%p4, %r14, 0;
+	@%p4 bra 	$L__BB0_4;
+	bra.uni 	$L__BB0_10;
+
+$L__BB0_4:
+	not.b32 	%r75, %r66;
+	add.s32 	%r76, %r13, %r75;
+	not.b32 	%r77, %r12;
+	max.s32 	%r78, %r76, %r77;
+	mov.u32 	%r79, -2;
+	sub.s32 	%r80, %r79, %r78;
+	and.b32  	%r240, %r14, 3;
+	setp.lt.u32 	%p5, %r80, 3;
+	mov.u32 	%r238, 0;
+	@%p5 bra 	$L__BB0_7;
+
+	sub.s32 	%r237, %r14, %r240;
+	mov.u32 	%r238, 0;
+
+$L__BB0_6:
+	add.s32 	%r82, %r238, %r13;
+	cvt.s64.s32 	%rd30, %r82;
+	add.s64 	%rd31, %rd1, %rd30;
+	ld.global.nc.u8 	%rs20, [%rd31];
+	add.s32 	%r83, %r8, %r238;
+	add.s32 	%r85, %r68, %r83;
+	st.volatile.shared.u8 	[%r85+1280], %rs20;
+	ld.global.nc.u8 	%rs21, [%rd31+1];
+	st.volatile.shared.u8 	[%r85+1281], %rs21;
+	ld.global.nc.u8 	%rs22, [%rd31+2];
+	st.volatile.shared.u8 	[%r85+1282], %rs22;
+	ld.global.nc.u8 	%rs23, [%rd31+3];
+	st.volatile.shared.u8 	[%r85+1283], %rs23;
+	add.s32 	%r238, %r238, 4;
+	add.s32 	%r237, %r237, -4;
+	setp.ne.s32 	%p6, %r237, 0;
+	@%p6 bra 	$L__BB0_6;
+
+$L__BB0_7:
+	setp.eq.s32 	%p7, %r240, 0;
+	@%p7 bra 	$L__BB0_10;
+
+	add.s32 	%r86, %r238, %r8;
+	add.s32 	%r88, %r68, %r86;
+	add.s32 	%r239, %r88, 1280;
+	add.s32 	%r89, %r238, %r13;
+	cvt.s64.s32 	%rd32, %r89;
+	add.s64 	%rd97, %rd1, %rd32;
+
+$L__BB0_9:
+	.pragma "nounroll";
+	ld.global.nc.u8 	%rs24, [%rd97];
+	st.volatile.shared.u8 	[%r239], %rs24;
+	add.s32 	%r239, %r239, 1;
+	add.s64 	%rd97, %rd97, 1;
+	add.s32 	%r240, %r240, -1;
+	setp.ne.s32 	%p8, %r240, 0;
+	@%p8 bra 	$L__BB0_9;
+
+$L__BB0_10:
+	add.s32 	%r27, %r3, 1284;
+	bar.sync 	0;
+	mul.lo.s32 	%r93, %r7, 5;
+	shr.s32 	%r94, %r93, 31;
+	shr.u32 	%r95, %r94, 29;
+	add.s32 	%r96, %r93, %r95;
+	shr.s32 	%r97, %r96, 3;
+	cvt.s64.s32 	%rd33, %r97;
+	cvta.to.global.u64 	%rd34, %rd24;
+	add.s64 	%rd35, %rd34, %rd33;
+	ld.global.nc.u8 	%rs25, [%rd35+1];
+	cvt.u32.u16 	%r98, %rs25;
+	and.b32  	%r99, %r98, 255;
+	ld.global.nc.u8 	%rs26, [%rd35];
+	cvt.u32.u16 	%r100, %rs26;
+	prmt.b32 	%r249, %r100, %r99, 30212;
+	mov.u32 	%r241, 0;
+	and.b32  	%r101, %r96, -8;
+	sub.s32 	%r102, %r101, %r93;
+	add.s32 	%r103, %r102, 11;
+	shr.u32 	%r104, %r249, %r103;
+	cvt.u64.u32 	%rd36, %r104;
+	cvt.u16.u32 	%rs27, %r104;
+	and.b16  	%rs45, %rs27, 31;
+	add.s32 	%r29, %r68, %r8;
+	ld.volatile.shared.u8 	%rd37, [%r29+1280];
+	shl.b64 	%rd38, %rd37, 56;
+	ld.volatile.shared.u8 	%rd39, [%r29+1281];
+	shl.b64 	%rd40, %rd39, 48;
+	or.b64  	%rd41, %rd40, %rd38;
+	ld.volatile.shared.u8 	%rd42, [%r29+1282];
+	shl.b64 	%rd43, %rd42, 40;
+	or.b64  	%rd44, %rd41, %rd43;
+	ld.volatile.shared.u8 	%rd45, [%r29+1283];
+	shl.b64 	%rd46, %rd45, 32;
+	or.b64  	%rd47, %rd44, %rd46;
+	ld.volatile.shared.u8 	%r105, [%r29+1284];
+	mul.wide.u32 	%rd48, %r105, 16777216;
+	or.b64  	%rd49, %rd47, %rd48;
+	ld.volatile.shared.u8 	%r106, [%r29+1285];
+	mul.wide.u32 	%rd50, %r106, 65536;
+	ld.volatile.shared.u8 	%r107, [%r29+1286];
+	mul.wide.u32 	%rd51, %r107, 256;
+	or.b64  	%rd52, %rd49, %rd50;
+	ld.volatile.shared.u8 	%rd53, [%r29+1287];
+	or.b64  	%rd54, %rd52, %rd51;
+	or.b64  	%rd55, %rd54, %rd53;
+	add.s32 	%r251, %r8, 8;
+	and.b64  	%rd5, %rd36, 31;
+	and.b32  	%r108, %r104, 31;
+	shl.b64 	%rd98, %rd55, %r108;
+	mov.u32 	%r109, -8;
+	sub.s32 	%r31, %r109, %r8;
+	cvta.to.global.u64 	%rd7, %rd25;
+	cvta.to.global.u64 	%rd8, %rd22;
+	cvta.to.global.u64 	%rd9, %rd23;
+	add.s32 	%r32, %r68, %r4;
+	mov.u16 	%rs43, %rs45;
+	mov.u32 	%r243, %r251;
+	bra.uni 	$L__BB0_11;
+
+$L__BB0_47:
+	add.s32 	%r224, %r68, %r243;
+	ld.volatile.shared.u8 	%r225, [%r224+1280];
+	shl.b32 	%r226, %r225, 24;
+	ld.volatile.shared.u8 	%r227, [%r224+1281];
+	shl.b32 	%r228, %r227, 16;
+	or.b32  	%r229, %r228, %r226;
+	ld.volatile.shared.u8 	%rs42, [%r224+1282];
+	mul.wide.u16 	%r230, %rs42, 256;
+	or.b32  	%r231, %r229, %r230;
+	ld.volatile.shared.u8 	%r232, [%r224+1283];
+	or.b32  	%r249, %r231, %r232;
+	add.s32 	%r243, %r243, 4;
+	cvt.u64.u32 	%rd92, %r249;
+	cvt.u64.u16 	%rd93, %rs43;
+	and.b64  	%rd94, %rd93, 255;
+	add.s64 	%rd95, %rd94, 4294967264;
+	cvt.u32.u64 	%r233, %rd95;
+	shl.b64 	%rd96, %rd92, %r233;
+	or.b64  	%rd98, %rd96, %rd98;
+	add.s16 	%rs43, %rs43, -32;
+	mov.u32 	%r241, %r36;
+
+$L__BB0_11:
+	shr.u64 	%rd56, %rd98, 56;
+	cvt.u32.u64 	%r110, %rd56;
+	add.s32 	%r112, %r68, %r110;
+	ld.volatile.shared.u8 	%rs44, [%r112];
+	setp.ne.s16 	%p9, %rs44, 0;
+	@%p9 bra 	$L__BB0_15;
+
+	shr.u64 	%rd57, %rd98, 48;
+	cvt.u32.u64 	%r113, %rd57;
+	and.b32  	%r114, %r113, 255;
+	add.s32 	%r116, %r68, %r114;
+	ld.volatile.shared.u8 	%rs44, [%r116+256];
+	setp.ne.s16 	%p10, %rs44, 0;
+	@%p10 bra 	$L__BB0_15;
+
+	shr.u64 	%rd58, %rd98, 40;
+	cvt.u32.u64 	%r117, %rd58;
+	and.b32  	%r118, %r117, 255;
+	add.s32 	%r120, %r68, %r118;
+	ld.volatile.shared.u8 	%rs44, [%r120+512];
+	setp.ne.s16 	%p11, %rs44, 0;
+	@%p11 bra 	$L__BB0_15;
+
+	shr.u64 	%rd59, %rd98, 32;
+	cvt.u32.u64 	%r121, %rd59;
+	and.b32  	%r122, %r121, 255;
+	add.s32 	%r124, %r68, %r122;
+	ld.volatile.shared.u8 	%rs44, [%r124+768];
+
+$L__BB0_15:
+	add.s32 	%r36, %r241, 1;
+	cvt.u32.u16 	%r125, %rs44;
+	and.b32  	%r126, %r125, 255;
+	add.s32 	%r128, %r68, %r126;
+	ld.volatile.shared.u8 	%rs28, [%r128+1024];
+	cvt.u32.u16 	%r129, %rs28;
+	and.b32  	%r130, %r129, 255;
+	shl.b64 	%rd98, %rd98, %r130;
+	add.s16 	%rs43, %rs28, %rs43;
+	and.b16  	%rs29, %rs43, 248;
+	shr.u16 	%rs30, %rs29, 3;
+	cvt.u32.u16 	%r131, %rs30;
+	add.s32 	%r132, %r31, %r243;
+	add.s32 	%r133, %r132, %r131;
+	setp.lt.u32 	%p12, %r133, %r65;
+	@%p12 bra 	$L__BB0_46;
+	bra.uni 	$L__BB0_16;
+
+$L__BB0_46:
+	and.b16  	%rs41, %rs43, 255;
+	setp.lt.u16 	%p32, %rs41, 32;
+	mov.u32 	%r241, %r36;
+	@%p32 bra 	$L__BB0_11;
+	bra.uni 	$L__BB0_47;
+
+$L__BB0_16:
+	setp.eq.s32 	%p13, %r252, 0;
+	shl.b32 	%r134, %r252, 2;
+	add.s32 	%r37, %r3, %r134;
+	add.s32 	%r38, %r3, %r69;
+	@%p13 bra 	$L__BB0_18;
+
+	add.s32 	%r234, %r241, 1;
+	st.volatile.shared.u32 	[%r37+1284], %r234;
+	bra.uni 	$L__BB0_19;
+
+$L__BB0_18:
+	mul.wide.u32 	%rd60, %r5, 4;
+	add.s64 	%rd61, %rd9, %rd60;
+	ld.global.nc.u32 	%r136, [%rd61];
+	st.volatile.shared.u32 	[%r38+1284], %r136;
+	ld.volatile.shared.u32 	%r137, [%r38+1284];
+	add.s32 	%r138, %r137, %r36;
+	st.volatile.shared.u32 	[%r3+1284], %r138;
+
+$L__BB0_19:
+	bar.sync 	0;
+	setp.lt.u32 	%p14, %r1, 2;
+	@%p14 bra 	$L__BB0_24;
+
+	add.s32 	%r39, %r252, 1;
+	mov.u32 	%r244, 2;
+
+$L__BB0_21:
+	rem.u32 	%r140, %r39, %r244;
+	setp.ne.s32 	%p15, %r140, 0;
+	@%p15 bra 	$L__BB0_23;
+
+	shr.u32 	%r141, %r244, 1;
+	sub.s32 	%r142, %r252, %r141;
+	shl.b32 	%r143, %r142, 2;
+	add.s32 	%r144, %r27, %r143;
+	ld.volatile.shared.u32 	%r145, [%r37+1284];
+	ld.volatile.shared.u32 	%r146, [%r144];
+	add.s32 	%r147, %r145, %r146;
+	st.volatile.shared.u32 	[%r37+1284], %r147;
+
+$L__BB0_23:
+	bar.sync 	0;
+	shl.b32 	%r244, %r244, 1;
+	setp.le.u32 	%p16, %r244, %r1;
+	@%p16 bra 	$L__BB0_21;
+
+$L__BB0_24:
+	setp.ne.s32 	%p17, %r252, 0;
+	@%p17 bra 	$L__BB0_26;
+
+	mov.u32 	%r148, 0;
+	st.volatile.shared.u32 	[%r38+1280], %r148;
+
+$L__BB0_26:
+	bar.sync 	0;
+	setp.lt.s32 	%p18, %r1, 2;
+	@%p18 bra 	$L__BB0_31;
+
+	add.s32 	%r42, %r252, 1;
+	mov.u32 	%r245, %r1;
+
+$L__BB0_28:
+	rem.u32 	%r149, %r42, %r245;
+	setp.eq.s32 	%p19, %r149, 0;
+	@%p19 bra 	$L__BB0_29;
+	bra.uni 	$L__BB0_30;
+
+$L__BB0_29:
+	shr.u32 	%r150, %r245, 1;
+	sub.s32 	%r151, %r252, %r150;
+	shl.b32 	%r152, %r151, 2;
+	add.s32 	%r153, %r27, %r152;
+	ld.volatile.shared.u32 	%r154, [%r37+1284];
+	ld.volatile.shared.u32 	%r155, [%r153];
+	add.s32 	%r156, %r154, %r155;
+	st.volatile.shared.u32 	[%r37+1284], %r156;
+	ld.volatile.shared.u32 	%r157, [%r153];
+	ld.volatile.shared.u32 	%r158, [%r37+1284];
+	sub.s32 	%r159, %r158, %r157;
+	st.volatile.shared.u32 	[%r153], %r159;
+
+$L__BB0_30:
+	shr.u32 	%r44, %r245, 1;
+	bar.sync 	0;
+	setp.gt.u32 	%p20, %r245, 3;
+	mov.u32 	%r245, %r44;
+	@%p20 bra 	$L__BB0_28;
+
+$L__BB0_31:
+	@%p3 bra 	$L__BB0_32;
+	bra.uni 	$L__BB0_33;
+
+$L__BB0_32:
+	ld.volatile.shared.u32 	%r160, [%r38+1284];
+	st.volatile.shared.u32 	[%r3+1284], %r160;
+	ld.volatile.shared.u32 	%r161, [%r37+1284];
+	add.s32 	%r162, %r161, %r36;
+	st.volatile.shared.u32 	[%r38+1284], %r162;
+
+$L__BB0_33:
+	bar.sync 	0;
+	ld.volatile.shared.u32 	%r163, [%r37+1284];
+	cvt.u64.u32 	%rd12, %r163;
+	ld.volatile.shared.u32 	%r45, [%r3+1284];
+	add.s32 	%r46, %r163, %r36;
+	ld.volatile.shared.u8 	%rd62, [%r29+1280];
+	shl.b64 	%rd63, %rd62, 56;
+	ld.volatile.shared.u8 	%rd64, [%r29+1281];
+	shl.b64 	%rd65, %rd64, 48;
+	or.b64  	%rd66, %rd65, %rd63;
+	ld.volatile.shared.u8 	%rd67, [%r29+1282];
+	shl.b64 	%rd68, %rd67, 40;
+	or.b64  	%rd69, %rd66, %rd68;
+	ld.volatile.shared.u8 	%rd70, [%r29+1283];
+	shl.b64 	%rd71, %rd70, 32;
+	or.b64  	%rd72, %rd69, %rd71;
+	ld.volatile.shared.u8 	%r164, [%r29+1284];
+	mul.wide.u32 	%rd73, %r164, 16777216;
+	ld.volatile.shared.u8 	%r165, [%r29+1285];
+	mul.wide.u32 	%rd74, %r165, 65536;
+	ld.volatile.shared.u8 	%r166, [%r29+1286];
+	mul.wide.u32 	%rd75, %r166, 256;
+	ld.volatile.shared.u8 	%rd76, [%r29+1287];
+	or.b64  	%rd77, %rd72, %rd73;
+	or.b64  	%rd78, %rd77, %rd74;
+	or.b64  	%rd79, %rd78, %rd75;
+	or.b64  	%rd80, %rd79, %rd76;
+	cvt.u32.u64 	%r167, %rd5;
+	shl.b64 	%rd100, %rd80, %r167;
+	setp.ge.u32 	%p22, %r163, %r67;
+	@%p22 bra 	$L__BB0_43;
+
+	cvt.u32.u64 	%r168, %rd12;
+	add.s64 	%rd99, %rd8, %rd12;
+	shl.b32 	%r169, %r168, 1;
+	add.s32 	%r170, %r4, %r169;
+	add.s32 	%r171, %r170, 1288;
+	shl.b32 	%r172, %r45, 1;
+	sub.s32 	%r173, %r171, %r172;
+	add.s32 	%r247, %r68, %r173;
+	add.s32 	%r246, %r168, 1;
+	bra.uni 	$L__BB0_35;
+
+$L__BB0_42:
+	add.s64 	%rd99, %rd99, 1;
+	add.s32 	%r247, %r247, 2;
+	add.s32 	%r246, %r246, 1;
+
+$L__BB0_35:
+	shr.u64 	%rd81, %rd100, 56;
+	cvt.u32.u64 	%r175, %rd81;
+	add.s32 	%r177, %r68, %r175;
+	ld.volatile.shared.u8 	%rs46, [%r177];
+	setp.ne.s16 	%p23, %rs46, 0;
+	@%p23 bra 	$L__BB0_39;
+
+	shr.u64 	%rd82, %rd100, 48;
+	cvt.u32.u64 	%r178, %rd82;
+	and.b32  	%r179, %r178, 255;
+	add.s32 	%r181, %r68, %r179;
+	ld.volatile.shared.u8 	%rs46, [%r181+256];
+	setp.ne.s16 	%p24, %rs46, 0;
+	@%p24 bra 	$L__BB0_39;
+
+	shr.u64 	%rd83, %rd100, 40;
+	cvt.u32.u64 	%r182, %rd83;
+	and.b32  	%r183, %r182, 255;
+	add.s32 	%r185, %r68, %r183;
+	ld.volatile.shared.u8 	%rs46, [%r185+512];
+	setp.ne.s16 	%p25, %rs46, 0;
+	@%p25 bra 	$L__BB0_39;
+
+	shr.u64 	%rd84, %rd100, 32;
+	cvt.u32.u64 	%r186, %rd84;
+	and.b32  	%r187, %r186, 255;
+	add.s32 	%r189, %r68, %r187;
+	ld.volatile.shared.u8 	%rs46, [%r189+768];
+
+$L__BB0_39:
+	ld.global.nc.u8 	%rs31, [%rd99];
+	and.b16  	%rs32, %rs31, 128;
+	and.b16  	%rs33, %rs46, 254;
+	shr.u16 	%rs34, %rs33, 1;
+	or.b16  	%rs35, %rs32, %rs34;
+	mul.wide.u16 	%r190, %rs35, 256;
+	and.b16  	%rs36, %rs31, 127;
+	cvt.u32.u16 	%r191, %rs46;
+	cvt.u32.u16 	%r192, %rs36;
+	bfi.b32 	%r193, %r191, %r192, 7, 9;
+	and.b32  	%r194, %r193, 255;
+	and.b32  	%r195, %r249, -65536;
+	or.b32  	%r196, %r195, %r194;
+	or.b32  	%r249, %r196, %r190;
+	st.volatile.shared.u16 	[%r247], %r249;
+	setp.ge.u32 	%p26, %r246, %r67;
+	setp.ge.u32 	%p27, %r246, %r46;
+	or.pred  	%p28, %p26, %p27;
+	@%p28 bra 	$L__BB0_43;
+
+	and.b32  	%r198, %r191, 255;
+	add.s32 	%r200, %r68, %r198;
+	ld.volatile.shared.u8 	%rs37, [%r200+1024];
+	cvt.u32.u16 	%r201, %rs37;
+	and.b32  	%r202, %r201, 255;
+	shl.b64 	%rd100, %rd100, %r202;
+	add.s16 	%rs45, %rs37, %rs45;
+	and.b16  	%rs38, %rs45, 255;
+	setp.lt.u16 	%p29, %rs38, 32;
+	@%p29 bra 	$L__BB0_42;
+
+	add.s32 	%r204, %r68, %r251;
+	ld.volatile.shared.u8 	%r205, [%r204+1280];
+	shl.b32 	%r206, %r205, 24;
+	ld.volatile.shared.u8 	%r207, [%r204+1281];
+	shl.b32 	%r208, %r207, 16;
+	or.b32  	%r209, %r208, %r206;
+	ld.volatile.shared.u8 	%rs39, [%r204+1282];
+	mul.wide.u16 	%r210, %rs39, 256;
+	or.b32  	%r211, %r209, %r210;
+	ld.volatile.shared.u8 	%r212, [%r204+1283];
+	or.b32  	%r249, %r211, %r212;
+	add.s32 	%r251, %r251, 4;
+	cvt.u64.u32 	%rd85, %r249;
+	cvt.u64.u16 	%rd86, %rs45;
+	and.b64  	%rd87, %rd86, 255;
+	add.s64 	%rd88, %rd87, 4294967264;
+	cvt.u32.u64 	%r213, %rd88;
+	shl.b64 	%rd89, %rd85, %r213;
+	or.b64  	%rd100, %rd89, %rd100;
+	add.s16 	%rs45, %rs45, -32;
+	bra.uni 	$L__BB0_42;
+
+$L__BB0_43:
+	bar.sync 	0;
+	ld.volatile.shared.u32 	%r214, [%r38+1284];
+	sub.s32 	%r215, %r214, %r45;
+	sub.s32 	%r60, %r67, %r45;
+	min.u32 	%r216, %r215, %r60;
+	setp.ge.u32 	%p30, %r252, %r216;
+	@%p30 bra 	$L__BB0_45;
+
+$L__BB0_44:
+	shl.b32 	%r217, %r252, 1;
+	add.s32 	%r218, %r32, %r217;
+	ld.volatile.shared.u16 	%rs40, [%r218+1288];
+	add.s32 	%r219, %r252, %r45;
+	mul.wide.u32 	%rd90, %r219, 2;
+	add.s64 	%rd91, %rd7, %rd90;
+	st.global.u16 	[%rd91], %rs40;
+	ld.volatile.shared.u32 	%r220, [%r38+1284];
+	sub.s32 	%r221, %r220, %r45;
+	min.u32 	%r222, %r221, %r60;
+	add.s32 	%r252, %r252, %r1;
+	setp.lt.u32 	%p31, %r252, %r222;
+	@%p31 bra 	$L__BB0_44;
+
+$L__BB0_45:
+	ret;
+
+}
+

layer_0.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:b4c01dcac753e1bca5c633fe7df1cd6d74a3cf425b7e3364778d9c20114074a5
+size 758603326

layer_1.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:294b42e51f93d30d3cc2cdeb8ed3a807a11dc775a3db6a8022374ccecc11d38a
+size 753471698

layer_10.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:509815b7173d47de56adb5f53e76ba790d44201e9ef52cada89ae3db152f5b68
+size 751783012

layer_11.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:909d2d5a6d13652b2243b8f1de932a07706db07a0b8313d497898785758c5baf
+size 751751314

layer_12.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:358282f3a74d48c83406fd32bcf314a4b903a962f0754b1dadca2f2622bfe466
+size 751875767

layer_13.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:6678559cd0e32f034fd57de4ab88251a9bde6336435b82e8d4824316895ab346
+size 751892664

layer_14.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:43e4616f6298bef9fbce486505a580772eb0d5eb0cd27ba52ed302c0c1238ee4
+size 751809969

layer_15.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:74d8660201d5a6e8639f2cb51a50789b02eecdca323ef865e73166c684ddb75b
+size 751676108

layer_16.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9e000973dd898635ddd8fac3e86458f2f22fd2a74d36aed2f67c5dc868651456
+size 751630430

layer_17.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:bdd7c745e2cacc59e9cd8d5de054009d417d8b9f955eaf5791d27898207ac74c
+size 751610977

layer_18.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:d67e93f6066e9bd30173505b4c656a3c6a0e45ea9389ad9a3a3a50e5db8681ce
+size 751474979

layer_19.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:e94ca47c23ee8050269a7abac3796d421c0c76e5f5bf6500813baad91e1effca
+size 751448233

layer_2.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:df9f492480b7ec3d8d305c7754b482ba1eb05c6c47fddfb9223cc1f0ad470c96
+size 754325380

layer_20.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:0f1df815fdecdd4ed84dbd02a3e95ad8fd525469a426ef466ddea33e7a4ce066
+size 751306616

layer_21.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:1ceb78dc40cbc3e7d1d0631558433f0789d23c89cf10d7ce7849ef7f0daaa9fd
+size 751282410

layer_22.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:12762755739b97e22179d5062857d8146d62c3888df15dc60df3a4e0d16f5308
+size 751293389

layer_23.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:3d4597269cd56615848163bc9ce7198abe444be464c27efcaffdf5269ab4265e
+size 751293812

layer_24.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:b4e1321b679c60332a334faf78cf1863f7b042507d85839cd95f7f37d8cb8eb1
+size 751193486

layer_25.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:297d65ce9acb406acb02bac73a475a3236de3f7df3c7752bcb742e7d2ad96d49
+size 751121450

layer_26.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:52a2689015f0b6da486fba1b804254627710244a22c1e9a2694be4deb7c4ddc8
+size 751052181

layer_27.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:ed31813d56c8c765a0395cfc39223675df2df2af86158eeab4889a147f3c0839
+size 751199805

layer_28.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:8d05913a7bd23f751580e949234fbf544389c256ea079f53597e030c680dafc3
+size 751120584

layer_29.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:1812d691b600ae6c9184500a00ba3ff1efc01aabc330a6f5b505732d7839afc8
+size 751093621

layer_3.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:8312b02f653fcd7897ac1f04a672eb89cc90609db319425099fe68984e48ac26
+size 752699382

layer_30.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:4716d07dd28fed5c5247c0a01e36a58fcb1f8e043d98433ae01938cbb17583de
+size 751033703

layer_31.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:a13a1624d363b857bae780cc92359a09887c1287ef46c178847f3851f553a216
+size 750950229

layer_32.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:2f72699b7e30a1ba233796182604123ee12de9056366ddd0730a94f279700584
+size 750969981

layer_33.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:cccf47a74e7b1875d71d19f3b8f3bb0c124b616f9145eb4572b03cc24a030775
+size 751016141

layer_34.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:74b120d0c47a012978423e49b2a5840f627e86e759d5e34fe87029f68102dc01
+size 751080560

layer_35.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:59025d4ade28ddc26bf39a3795d761e4997bc58481be9e0d451d2a0c959a2713
+size 751199796

layer_36.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:99dbfff5b2c742adeca6d759fa3c1636047bcc329545dbf4c319dad949b47800
+size 751278784

layer_37.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:cdbf4cc5eb186ab18cd97569cd5475e17189d296f0732330ed4817214d69e77d
+size 751610563

layer_38.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:6a6b3e620b241cfb278e2e827b38972f396bacd49e795c51b2ee2cd2fab80617
+size 751681420

layer_39.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:4cef6e2c8fcbbf71868d5de86279b4a766351a8f2b5f8108626be35d3b5ff851
+size 753153030

layer_4.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:b9b195a6cf53c7b650c17662a75647d0f519a5000d470a4f3d1998aae4a63b47
+size 752543953

layer_5.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:7ef0189c5d2d781bc6e4865daf8dca279332c472ad29e2105ea6fbe410cfa27d
+size 752135655

layer_6.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:97e4345c015c3dbe51a26a66b6ce2b7fba1adb4dca87c3c2fdbd6558dd8af736
+size 752335495

layer_7.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:2f2b15310a708dcc585c123b662273047a308007e052b2e9666cbe65bed10181
+size 752082894

layer_8.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:4a31472c2e5e2316d13c7cab78db48b19d2116ad197d176e2abee8462378e802
+size 751955327

layer_9.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:35d8da267336d290326e310f2024ea0757ca72989e58368cef97f40be9701388
+size 751822822

lm_head.pkl

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9b0cd341e6c6a7ed0d21a6e8d680232c6aa35d901a43437604014c14e6fe5cec
+size 910765294