| // |
| // 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; |
|
|
| } |
|
|
|
|
