From e177edcdefca06d7d3a7bc7be5f6a759f69b909e Mon Sep 17 00:00:00 2001 From: aamir Date: Fri, 1 Jun 2018 11:18:17 -0700 Subject: remove unwanted files --- cuda-kernels/_cuobjdump_1.elf | 15 - cuda-kernels/_cuobjdump_1.ptx | 170 - cuda-kernels/_cuobjdump_1.sass | 2 - cuda-kernels/_cuobjdump_2.elf | 494 -- cuda-kernels/_cuobjdump_2.sass | 348 -- cuda-kernels/_cuobjdump_complete_output_EIGzTK | 1055 ---- cuda-kernels/_cuobjdump_complete_output_rndQyq | 1055 ---- cuda-kernels/gpgpu_inst_stats.txt | 20 - ...usim_power_report__Sun-May-27-14-17-34-2018.log | 324 - ...usim_power_report__Sun-May-27-14-17-47-2018.log | 324 - cuda-kernels/log | 6328 -------------------- cuda-kernels/log1 | 512 -- cuda-kernels/tensor_core | Bin 2750968 -> 0 bytes 13 files changed, 10647 deletions(-) delete mode 100644 cuda-kernels/_cuobjdump_1.elf delete mode 100644 cuda-kernels/_cuobjdump_1.ptx delete mode 100644 cuda-kernels/_cuobjdump_1.sass delete mode 100644 cuda-kernels/_cuobjdump_2.elf delete mode 100644 cuda-kernels/_cuobjdump_2.sass delete mode 100644 cuda-kernels/_cuobjdump_complete_output_EIGzTK delete mode 100644 cuda-kernels/_cuobjdump_complete_output_rndQyq delete mode 100755 cuda-kernels/gpgpu_inst_stats.txt delete mode 100644 cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-34-2018.log delete mode 100644 cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-47-2018.log delete mode 100644 cuda-kernels/log delete mode 100644 cuda-kernels/log1 delete mode 100755 cuda-kernels/tensor_core diff --git a/cuda-kernels/_cuobjdump_1.elf b/cuda-kernels/_cuobjdump_1.elf deleted file mode 100644 index 672b0f0..0000000 --- a/cuda-kernels/_cuobjdump_1.elf +++ /dev/null @@ -1,15 +0,0 @@ -64bit elf: type=2, abi=7, sm=70, toolkit=90, flags = 0x460546 -Sections: -Index Offset Size ES Align Type Flags Link Info Name - 1 40 32 0 1 STRTAB 0 0 0 .shstrtab - 2 72 32 0 1 STRTAB 0 0 0 .strtab - 3 a8 18 18 8 SYMTAB 0 2 0 .symtab - -.section .strtab - -.section .shstrtab - -.section .symtab - index value size info other shndx name - 0 0 0 0 0 0 (null) - diff --git a/cuda-kernels/_cuobjdump_1.ptx b/cuda-kernels/_cuobjdump_1.ptx deleted file mode 100644 index 3453f4a..0000000 --- a/cuda-kernels/_cuobjdump_1.ptx +++ /dev/null @@ -1,170 +0,0 @@ - - - - - - - -.version 6.0 -.target sm_70 -.address_size 64 - - -.extern .func (.param .b32 func_retval0) vprintf -( -.param .b64 vprintf_param_0, -.param .b64 vprintf_param_1 -) -; -.global .align 16 .b8 $str[9] = {99, 108, 111, 99, 107, 61, 37, 100, 0}; - -.visible .entry _Z12wmma_exampleP6__halfS0_Pfiiiff( -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_0, -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_1, -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_2, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_3, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_4, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_5, -.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_6, -.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_7 -) -{ -.local .align 8 .b8 __local_depot0[8]; -.reg .b64 %SP; -.reg .b64 %SPL; -.reg .pred %p<6>; -.reg .f32 %f<34>; -.reg .b32 %r<38>; -.reg .b64 %rd<18>; - - -mov.u64 %rd17, __local_depot0; -cvta.local.u64 %SP, %rd17; -ld.param.u64 %rd1, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0]; -ld.param.u64 %rd2, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1]; -ld.param.u64 %rd3, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2]; -ld.param.u32 %r4, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3]; -ld.param.u32 %r7, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4]; -ld.param.u32 %r5, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5]; - - mov.u32 %r6, %clock; - - mov.u32 %r8, %ntid.x; -mov.u32 %r9, %ctaid.x; -mov.u32 %r10, %tid.x; -mad.lo.s32 %r11, %r8, %r9, %r10; -mov.u32 %r12, WARP_SZ; -div.u32 %r13, %r11, %r12; -mov.u32 %r14, %ntid.y; -mov.u32 %r15, %ctaid.y; -mov.u32 %r16, %tid.y; -mad.lo.s32 %r17, %r14, %r15, %r16; -shl.b32 %r2, %r13, 4; -shl.b32 %r3, %r17, 4; -setp.lt.s32 %p1, %r2, %r4; -setp.gt.s32 %p2, %r5, 0; -and.pred %p3, %p1, %p2; -setp.lt.s32 %p4, %r3, %r7; -and.pred %p5, %p3, %p4; -mov.f32 %f26, 0f00000000; -mov.f32 %f27, %f26; -mov.f32 %f28, %f26; -mov.f32 %f29, %f26; -mov.f32 %f30, %f26; -mov.f32 %f31, %f26; -mov.f32 %f32, %f26; -mov.f32 %f33, %f26; -@!%p5 bra BB0_2; -bra.uni BB0_1; - -BB0_1: -mul.wide.s32 %rd4, %r2, 2; -add.s64 %rd5, %rd1, %rd4; -wmma.load.a.sync.row.m16n16k16.f16 {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, [%rd5], %r4; -mul.wide.s32 %rd6, %r3, 2; -add.s64 %rd7, %rd2, %rd6; -wmma.load.b.sync.col.m16n16k16.f16 {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, [%rd7], %r5; -mov.f32 %f25, 0f00000000; -wmma.mma.sync.row.col.m16n16k16.f32.f32 {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, {%f25, %f25, %f25, %f25, %f25, %f25, %f25, %f25}; - -BB0_2: -add.u64 %rd8, %SP, 0; -cvta.to.local.u64 %rd9, %rd8; -mul.lo.s32 %r35, %r3, %r4; -cvt.s64.s32 %rd10, %r35; -cvt.s64.s32 %rd11, %r2; -add.s64 %rd12, %rd10, %rd11; -shl.b64 %rd13, %rd12, 2; -add.s64 %rd14, %rd3, %rd13; -wmma.store.d.sync.col.m16n16k16.f32 [%rd14], {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, %r4; - - mov.u32 %r34, %clock; - - sub.s32 %r36, %r34, %r6; -st.local.u32 [%rd9], %r36; -mov.u64 %rd15, $str; -cvta.global.u64 %rd16, %rd15; - - { -.reg .b32 temp_param_reg; - - .param .b64 param0; -st.param.b64 [param0+0], %rd16; -.param .b64 param1; -st.param.b64 [param1+0], %rd8; -.param .b32 retval0; -call.uni (retval0), -vprintf, -( -param0, -param1 -); -ld.param.b32 %r37, [retval0+0]; - - - } - ret; -} - - -.visible .entry _Z17convertFp32ToFp16P6__halfPfi( -.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_0, -.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_1, -.param .u32 _Z17convertFp32ToFp16P6__halfPfi_param_2 -) -{ -.reg .pred %p<2>; -.reg .b16 %rs<2>; -.reg .f32 %f<2>; -.reg .b32 %r<6>; -.reg .b64 %rd<9>; - - -ld.param.u64 %rd1, [_Z17convertFp32ToFp16P6__halfPfi_param_0]; -ld.param.u64 %rd2, [_Z17convertFp32ToFp16P6__halfPfi_param_1]; -ld.param.u32 %r2, [_Z17convertFp32ToFp16P6__halfPfi_param_2]; -mov.u32 %r3, %ntid.x; -mov.u32 %r4, %ctaid.x; -mov.u32 %r5, %tid.x; -mad.lo.s32 %r1, %r4, %r3, %r5; -setp.ge.s32 %p1, %r1, %r2; -@%p1 bra BB1_2; - -cvta.to.global.u64 %rd3, %rd2; -mul.wide.s32 %rd4, %r1, 4; -add.s64 %rd5, %rd3, %rd4; -ld.global.f32 %f1, [%rd5]; - - { cvt.rn.f16.f32 %rs1, %f1;} - - - cvta.to.global.u64 %rd6, %rd1; -mul.wide.s32 %rd7, %r1, 2; -add.s64 %rd8, %rd6, %rd7; -st.global.u16 [%rd8], %rs1; - -BB1_2: -ret; -} - - diff --git a/cuda-kernels/_cuobjdump_1.sass b/cuda-kernels/_cuobjdump_1.sass deleted file mode 100644 index 2aac29a..0000000 --- a/cuda-kernels/_cuobjdump_1.sass +++ /dev/null @@ -1,2 +0,0 @@ - code for sm_70 - diff --git a/cuda-kernels/_cuobjdump_2.elf b/cuda-kernels/_cuobjdump_2.elf deleted file mode 100644 index c03b06d..0000000 --- a/cuda-kernels/_cuobjdump_2.elf +++ /dev/null @@ -1,494 +0,0 @@ -64bit elf: type=2, abi=7, sm=70, toolkit=90, flags = 0x460546 -Sections: -Index Offset Size ES Align Type Flags Link Info Name - 1 40 21b 0 1 STRTAB 0 0 0 .shstrtab - 2 25b 273 0 1 STRTAB 0 0 0 .strtab - 3 4d0 108 18 8 SYMTAB 0 2 7 .symtab - 4 5d8 e0 0 1 PROGBITS 0 0 0 .debug_frame - 5 6b8 48 0 4 CUDA_INFO 0 3 0 .nv.info - 6 700 50 0 4 CUDA_INFO 0 3 d .nv.info._Z17convertFp32ToFp16P6__halfPfi - 7 750 ac 0 4 CUDA_INFO 0 3 e .nv.info._Z12wmma_exampleP6__halfS0_Pfiiiff - 8 800 30 10 8 REL 0 3 e .rel.text._Z12wmma_exampleP6__halfS0_Pfiiiff - 9 830 30 18 8 RELA 0 3 e .rela.text._Z12wmma_exampleP6__halfS0_Pfiiiff - a 860 20 10 8 REL 0 3 4 .rel.debug_frame - b 880 174 0 4 PROGBITS 2 0 d .nv.constant0._Z17convertFp32ToFp16P6__halfPfi - c 9f4 18c 0 4 PROGBITS 2 0 e .nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff - d b80 100 0 80 PROGBITS 6 3 9000008 .text._Z17convertFp32ToFp16P6__halfPfi - e c80 980 0 80 PROGBITS 6 3 20000009 .text._Z12wmma_exampleP6__halfS0_Pfiiiff - f 1600 9 0 10 PROGBITS 3 0 0 .nv.global.init - -.section .strtab - -.section .shstrtab - -.section .symtab - index value size info other shndx name - 0 0 0 0 0 0 (null) - 1 0 0 3 0 d .text._Z17convertFp32ToFp16P6__halfPfi - 2 0 0 3 0 f .nv.global.init - 3 0 9 1 0 f $str - 4 0 0 3 0 b .nv.constant0._Z17convertFp32ToFp16P6__halfPfi - 5 0 0 3 0 e .text._Z12wmma_exampleP6__halfS0_Pfiiiff - 6 0 0 3 0 c .nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff - 7 0 0 3 0 4 .debug_frame - 8 0 256 12 10 d _Z17convertFp32ToFp16P6__halfPfi - 9 0 2432 12 10 e _Z12wmma_exampleP6__halfS0_Pfiiiff - 10 0 0 12 0 0 vprintf - - -.nv.constant0._Z17convertFp32ToFp16P6__halfPfi -0x00000000 0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 - - - -.nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff -0x00000000 0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 - - -.nv.global.init -0x636f6c63 0x64253d6b 0 - - -.nv.info - <0x1> - Attribute: EIATTR_MAX_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x9 0x0 - <0x2> - Attribute: EIATTR_MIN_STACK_SIZE - Format: EIFMT_SVAL - Value: function: _Z12wmma_exampleP6__halfS0_Pfiiiff(0x9) min stack size: 0x8 - <0x3> - Attribute: EIATTR_FRAME_SIZE - Format: EIFMT_SVAL - Value: function: _Z12wmma_exampleP6__halfS0_Pfiiiff(0x9) frame size: 0x8 - <0x4> - Attribute: EIATTR_MAX_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x8 0x0 - <0x5> - Attribute: EIATTR_MIN_STACK_SIZE - Format: EIFMT_SVAL - Value: function: _Z17convertFp32ToFp16P6__halfPfi(0x8) min stack size: 0x0 - <0x6> - Attribute: EIATTR_FRAME_SIZE - Format: EIFMT_SVAL - Value: function: _Z17convertFp32ToFp16P6__halfPfi(0x8) frame size: 0x0 - - -.nv.info._Z17convertFp32ToFp16P6__halfPfi - <0x1> - Attribute: EIATTR_PARAM_CBANK - Format: EIFMT_SVAL - Value: 0x4 0x140160 - <0x2> - Attribute: EIATTR_CBANK_PARAM_SIZE - Format: EIFMT_HVAL - Value: 0x14 - <0x3> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x2 Offset : 0x10 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x4> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x1 Offset : 0x8 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x5> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x0 Offset : 0x0 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x6> - Attribute: EIATTR_MAXREG_COUNT - Format: EIFMT_HVAL - Value: 0xff - <0x7> - Attribute: EIATTR_EXIT_INSTR_OFFSETS - Format: EIFMT_SVAL - Value: 0x60 0xe0 - - -.nv.info._Z12wmma_exampleP6__halfS0_Pfiiiff - <0x1> - Attribute: EIATTR_PARAM_CBANK - Format: EIFMT_SVAL - Value: 0x6 0x2c0160 - <0x2> - Attribute: EIATTR_CBANK_PARAM_SIZE - Format: EIFMT_HVAL - Value: 0x2c - <0x3> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x7 Offset : 0x28 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x4> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x6 Offset : 0x24 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x5> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x5 Offset : 0x20 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x6> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x4 Offset : 0x1c Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x7> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x3 Offset : 0x18 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x8> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x2 Offset : 0x10 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x9> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x1 Offset : 0x8 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x10> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x0 Offset : 0x0 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x11> - Attribute: EIATTR_MAXREG_COUNT - Format: EIFMT_HVAL - Value: 0xff - <0x12> - Attribute: EIATTR_EXIT_INSTR_OFFSETS - Format: EIFMT_SVAL - Value: 0x940 - <0x13> - Attribute: EIATTR_EXTERNS - Format: EIFMT_SVAL - Value: externs: vprintf(0xa) - <0x14> - Attribute: EIATTR_CRS_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x0 - - -.text._Z17convertFp32ToFp16P6__halfPfi -bar = 0 reg = 9 lmem=0 smem=0 -0xfffff389 0x000000ff 0x000e00ff 0x000fe200 -0x00017a02 0x00000a00 0x00000f00 0x000fd000 -0x00047919 0x00000000 0x00002500 0x000e2200 -0x00027919 0x00000000 0x00002100 0x000e2400 -0x04047a24 0x00000000 0x078e0202 0x001fca00 -0x04007a0c 0x00005c00 0x03f062f0 0x000fd800 -0x0000094d 0x00000000 0x03800000 0x000fea00 -0x00027802 0x00000004 0x00000f00 0x000fca00 -0x04027625 0x00005a00 0x078e0202 0x000fd400 -0x02027381 0x00000000 0x001ee900 0x000e2200 -0x00057802 0x00000002 0x00000f00 0x000fca00 -0x04047625 0x00005800 0x078e0205 0x000fe200 -0x00067304 0x00000002 0x00200800 0x001e3200 -0x04007386 0x00000006 0x0010e500 0x0011e200 -0x0000794d 0x00000000 0x03800000 0x000fea00 -0x00007947 0xfffffff0 0x0383ffff 0x000fc000 - - - -.text._Z12wmma_exampleP6__halfS0_Pfiiiff -bar = 0 reg = 32 lmem=0 smem=0 -0xfffff389 0x000000ff 0x000e00ff 0x000fe200 -0xff017624 0x00000a00 0x078e00ff 0x000fd000 -0x01017810 0xfffffff8 0x07ffe0ff 0x000fc800 -0x01027a10 0x00000800 0x07f1e0ff 0x000fca00 -0xff007624 0x00000900 0x000e06ff 0x000fd000 -0x00037805 0x00000000 0x00005000 0x000fd000 -0x00077906 0x00000020 0x00209000 0x000e2400 -0x00077308 0x00000007 0x00001000 0x001e2200 -0x00067919 0x00000000 0x00002500 0x000e6200 -0x00097919 0x00000000 0x00002100 0x000e6200 -0x07087810 0x0ffffffe 0x07ffe0ff 0x001fcc00 -0x00057305 0x00000008 0x0021f000 0x0000a200 -0xff047224 0x000000ff 0x078e00ff 0x000fe400 -0x06067a24 0x00000000 0x078e0209 0x002fe400 -0x050a7824 0xffffffe0 0x078e00ff 0x004fc800 -0x05047225 0x0000000a 0x078e0004 0x000fd000 -0x05047225 0x00000006 0x078e00ff 0x000fcc00 -0xff047224 0x000000ff 0x078e0a05 0x000fc800 -0x04067824 0x00000020 0x078e0206 0x000fca00 -0x0600780c 0x00000020 0x03f060f0 0x040fe200 -0x001c7919 0x00000000 0x00002600 0x000e2200 -0x00077919 0x00000000 0x00002200 0x000e3400 -0x06060810 0xffffffe0 0x07ffe0ff 0x000fc800 -0x0600780c 0x00000020 0x03f260f0 0x000fe400 -0x05050810 0x00000001 0x07ffe0ff 0x000fe400 -0xff007a0c 0x00006000 0x03f012f0 0x000fd000 -0x05051810 0x00000001 0x07ffe0ff 0x000fe200 -0x1c1c7a24 0x00000100 0x078e0207 0x001fc600 -0x051d7819 0x00000004 0x000006ff 0x000fe200 -0x1c1c7824 0x00000010 0x078e00ff 0x000fc600 -0x1d007a0c 0x00005e00 0x007012f0 0x000fc800 -0x1c007a0c 0x00005f00 0x007012f0 0x000fe200 -0x00007945 0x000003a0 0x03800000 0x000fe200 -0xff077224 0x000000ff 0x078e00ff 0x000fe200 -0x000b7202 0x000000ff 0x00000f00 0x000fe200 -0xff067224 0x000000ff 0x078e00ff 0x000fe400 -0xff057224 0x000000ff 0x078e00ff 0x000fe400 -0xff047224 0x000000ff 0x078e00ff 0x000fe400 -0xff0a7224 0x000000ff 0x078e00ff 0x000fc400 -0xff097224 0x000000ff 0x078e00ff 0x000fe400 -0xff087224 0x000000ff 0x078e00ff 0x000fe200 -0x00008947 0x00000300 0x03800000 0x000fee00 -0x00067919 0x00000000 0x00000000 0x000e2200 -0xff0a7424 0x00000002 0x078e00ff 0x000fc800 -0x1d107625 0x00005800 0x078e020a 0x000fe200 -0xff047819 0x00000002 0x00011606 0x001fc800 -0x04057812 0x00000003 0x078ec0ff 0x000fe400 -0x06047812 0x00000003 0x078ec0ff 0x000fe400 -0x05077812 0x00000001 0x078ec0ff 0x000fe400 -0xff067819 0x00000004 0x00011606 0x000fe400 -0xff057819 0x00000001 0x00011605 0x000fe200 -0x07077824 0x00000008 0x078e0204 0x000fe200 -0x06067812 0x00000001 0x078ec0ff 0x000fc400 -0x05047211 0x00000004 0x078e18ff 0x000fe200 -0x1c0c7625 0x00005a00 0x078e020a 0x000fe400 -0x06077824 0x00000004 0x078e0207 0x040fe400 -0x06047824 0x00000004 0x078e0204 0x000fe400 -0x07077824 0x00000002 0x078e00ff 0x000fe400 -0x04057824 0x00000002 0x078e00ff 0x000fe400 -0x07107a25 0x00005e00 0x078e0010 0x000fc400 -0x050c7a25 0x00006000 0x078e000c 0x000fd000 -0x10187980 0x00000000 0x0010ed00 0x00006400 -0x0c147980 0x00000000 0x0010ed00 0x00046200 -0x10107980 0x00000010 0x0010ed00 0x001e2200 -0x0c0c7980 0x00000010 0x0010ed00 0x004e2200 -0xff087224 0x000000ff 0x078e00ff 0x000fe200 -0x00097202 0x000000ff 0x00000f00 0x000fe200 -0xff0a7224 0x000000ff 0x078e00ff 0x000fe400 -0xff0b7224 0x000000ff 0x078e00ff 0x000fe200 -0x00077202 0x000000ff 0x00000f00 0x000fe200 -0xff047224 0x000000ff 0x078e00ff 0x000fc400 -0xff057224 0x000000ff 0x078e00ff 0x000fe400 -0xff067224 0x000000ff 0x078e00ff 0x000fe200 -0x00007948 0xffffffff 0x03800000 0x000fe200 -0x18087236 0x00000014 0x00005408 0x0c226400 -0x180a7236 0x00000014 0x0000d40a 0x0c04a400 -0x18047236 0x00000014 0x00015404 0x0c06e400 -0x18067236 0x00000014 0x0001d406 0x00092800 -0x1a087236 0x00000016 0x00005408 0x0c202400 -0x1a0a7236 0x00000016 0x0000d40a 0x0c426400 -0x1a047236 0x00000016 0x00015404 0x0c84a400 -0x1a067236 0x00000016 0x0001d406 0x0106e800 -0x10087236 0x0000000c 0x00005408 0x0c102400 -0x100a7236 0x0000000c 0x0000d40a 0x0c226400 -0x10047236 0x0000000c 0x00015404 0x0c44a400 -0x10067236 0x0000000c 0x0001d406 0x0086e800 -0x12087236 0x0000000e 0x00005408 0x0c102400 -0x120a7236 0x0000000e 0x0000d40a 0x0c202400 -0x12047236 0x0000000e 0x00015404 0x0c402400 -0x12067236 0x0000000e 0x0001d406 0x00803400 -0x00007941 0x00000000 0x03800000 0x001fea00 -0x000c7919 0x00000000 0x00000000 0x000e2200 -0x1c1c7a24 0x00005e00 0x078e02ff 0x000fe200 -0xff0e7819 0x00000004 0x0001160c 0x001fc400 -0xff0d7819 0x00000002 0x0001160c 0x000fe400 -0x0c0c7812 0x00000003 0x078ec0ff 0x000fe400 -0x0e0e7812 0x00000001 0x078ec0ff 0x000fe400 -0x0d0d7812 0x00000003 0x078ec0ff 0x000fc600 -0x0e0c7824 0x00000004 0x078e020c 0x000fe200 -0x0d0f7812 0x00000001 0x078ec0ff 0x000fe400 -0xff107819 0x00000001 0x0001160d 0x000fe400 -0x0c0d7812 0x00000005 0x078ec0ff 0x040fe400 -0x0c0e7812 0x00000002 0x078ec0ff 0x000fc600 -0x0f0c7824 0x00000008 0x078e020d 0x000fe200 -0xff0f7819 0x0000001f 0x0001141d 0x000fe200 -0x100e7824 0x00000008 0x078e020e 0x000fe200 -0x1d117210 0x0000001c 0x07f1e0ff 0x000fe200 -0xff0d7224 0x000000ff 0x078e00ff 0x000fc600 -0x1c0f7211 0x0000000f 0x000f0eff 0x000fe200 -0x0e0c7a25 0x00005e00 0x078e000c 0x000fe200 -0x110e7a11 0x00005c00 0x078010ff 0x000fe200 -0xff127624 0x00005e00 0x078e00ff 0x000fc600 -0x11117a11 0x00005d00 0x000f140f 0x000fe400 -0x0c107211 0x0000000e 0x078010ff 0x000fe400 -0x120e7819 0x00000002 0x000006ff 0x000fe400 -0xff0f7819 0x0000001e 0x00011612 0x000fe400 -0x0c0d7211 0x00000011 0x000f140d 0x000fe400 -0x0e127211 0x00000010 0x078210ff 0x000fc400 -0x0e117210 0x00000010 0x07f1e0ff 0x040fe400 -0x0e137211 0x0000000d 0x008f140f 0x040fe400 -0x0e157210 0x00000012 0x07f3e0ff 0x000fe200 -0x0f147824 0x00000001 0x000e060d 0x040fe400 -0xff0c7224 0x000000ff 0x078e0010 0x000fe400 -0x0f167824 0x00000001 0x008e0613 0x000fe400 -0xff0e7224 0x000000ff 0x078e0011 0x000fc400 -0xff0f7224 0x000000ff 0x078e0014 0x000fe200 -0x00107202 0x00000015 0x00000f00 0x000fe200 -0xff117224 0x000000ff 0x078e0016 0x000fe200 -0x0c007385 0x00000000 0x0010e908 0x0001e200 -0x0c007385 0x00000008 0x0010e90a 0x0003e800 -0x0e007385 0x00000000 0x0010e909 0x0003e200 -0x0e007385 0x00000008 0x0010e90b 0x0003e200 -0x12007385 0x00000000 0x0010e904 0x0003e200 -0x12007385 0x00000008 0x0010e906 0x0003e200 -0x10007385 0x00000000 0x0010e905 0x0003e200 -0x10007385 0x00000008 0x0010e907 0x0003e200 -0x00007948 0xffffffff 0x03800000 0x000fe200 -0x02087a10 0x80000800 0x07ffe0ff 0x001fd000 -0x00047805 0x00000000 0x00005000 0x002fd000 -0x04037824 0x00000001 0x078e0a03 0x000fd000 -0x08007387 0x00000003 0x00100800 0x0001e200 -0xff067224 0x000000ff 0x078e0002 0x000fe200 -0x00047802 0x00000000 0x00000f00 0x000fe200 -0xff077224 0x000000ff 0x078e0000 0x000fe200 -0x00057802 0x00000000 0x00000f00 0x000fe400 -0x00147802 0x00000000 0x00000f00 0x000fe400 -0x00157802 0x00000000 0x00000f00 0x000fd000 -0x00007943 0x00000000 0x03c00000 0x001fea00 -0x0000794d 0x00000000 0x03800000 0x000fea00 -0x00007947 0xfffffff0 0x0383ffff 0x000fc000 -0x00007918 0x00000000 0x00000000 0x000fc000 -0x00007918 0x00000000 0x00000000 0x000fc000 - - -.section .rel.text._Z12wmma_exampleP6__halfS0_Pfiiiff REL -2272 $str R_CUDA_ABS32_LO_32 -2304 $str R_CUDA_ABS32_HI_32 -2352 vprintf R_CUDA_ABS47_34 - -.section .rela.text._Z12wmma_exampleP6__halfS0_Pfiiiff RELA -2320 _Z12wmma_exampleP6__halfS0_Pfiiiff R_CUDA_ABS32_LO_32 2368 -2336 _Z12wmma_exampleP6__halfS0_Pfiiiff R_CUDA_ABS32_HI_32 2368 - -.section .debug_frame -decodeDebugFrame, frameBuf 0xffffffff, total_length 224 -CIE length 40, cie_id -1 -version 3 -augmentation slen 1 -augmentation -code_align_factor slen 1 -data_align_factor slen 1 - Debug Frame Common Information Entry - length: 40 - CIE_id : -1 - version: 3 - augmentation: - code align factor: 4 - data align factor: -4 - return address register 16777215 - initial instructions: 23 bytes, ptr = 0x8080810c, frameBuf = 0xffffffff - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_same_value R255 - DW_CFA_same_value R1 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - Debug Frame Description Entry - length: 48 - CIE_pointer: 0 - initial_location: 0x0 - address_range: 0x100 - instructions: 24 bytes - DW_CFA_advance_loc4 delta 4 - DW_CFA_advance_loc4 delta 0 - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_advance_loc4 delta 52 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop -CIE length 40, cie_id -1 -version 3 -augmentation slen 1 -augmentation -code_align_factor slen 1 -data_align_factor slen 1 - Debug Frame Common Information Entry - length: 40 - CIE_id : -1 - version: 3 - augmentation: - code align factor: 4 - data align factor: -4 - return address register 16777215 - initial instructions: 23 bytes, ptr = 0x8080810c, frameBuf = 0xffffffff - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_same_value R255 - DW_CFA_same_value R1 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - Debug Frame Description Entry - length: 48 - CIE_pointer: 0 - initial_location: 0x0 - address_range: 0x970 - instructions: 24 bytes - DW_CFA_advance_loc4 delta 4 - DW_CFA_advance_loc4 delta 2 - DW_CFA_def_cfa register R1, offset 8 - DW_CFA_advance_loc4 delta 586 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - -.section .rel.debug_frame REL -72 _Z17convertFp32ToFp16P6__halfPfi R_NV_64 -184 _Z12wmma_exampleP6__halfS0_Pfiiiff R_NV_64 - diff --git a/cuda-kernels/_cuobjdump_2.sass b/cuda-kernels/_cuobjdump_2.sass deleted file mode 100644 index 1b50ed2..0000000 --- a/cuda-kernels/_cuobjdump_2.sass +++ /dev/null @@ -1,348 +0,0 @@ - code for sm_70 - Function : _Z17convertFp32ToFp16P6__halfPfi - .headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)" - /*0000*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ; /* 0x000000fffffff389 */ - /* 0x000fe200000e00ff */ - /*0010*/ MOV R1, c[0x0][0x28]; /* 0x00000a0000017a02 */ - /* 0x000fd00000000f00 */ - /*0020*/ S2R R4, SR_CTAID.X; /* 0x0000000000047919 */ - /* 0x000e220000002500 */ - /*0030*/ S2R R2, SR_TID.X; /* 0x0000000000027919 */ - /* 0x000e240000002100 */ - /*0040*/ IMAD R4, R4, c[0x0][0x0], R2; /* 0x0000000004047a24 */ - /* 0x001fca00078e0202 */ - /*0050*/ ISETP.GE.AND P0, PT, R4, c[0x0][0x170], PT, !PT; /* 0x00005c0004007a0c */ - /* 0x000fd80003f062f0 */ - /*0060*/ @P0 EXIT; /* 0x000000000000094d */ - /* 0x000fea0003800000 */ - /*0070*/ MOV R2, 0x4; /* 0x0000000400027802 */ - /* 0x000fca0000000f00 */ - /*0080*/ IMAD.WIDE R2, R4, R2, c[0x0][0x168]; /* 0x00005a0004027625 */ - /* 0x000fd400078e0202 */ - /*0090*/ LDG.E.SYS R2, [R2]; /* 0x0000000002027381 */ - /* 0x000e2200001ee900 */ - /*00a0*/ MOV R5, 0x2; /* 0x0000000200057802 */ - /* 0x000fca0000000f00 */ - /*00b0*/ IMAD.WIDE R4, R4, R5, c[0x0][0x160]; /* 0x0000580004047625 */ - /* 0x000fe200078e0205 */ - /*00c0*/ F2F.F16.F32 R6, R2; /* 0x0000000200067304 */ - /* 0x001e320000200800 */ - /*00d0*/ STG.E.U16.SYS [R4], R6; /* 0x0000000604007386 */ - /* 0x0011e2000010e500 */ - /*00e0*/ EXIT; /* 0x000000000000794d */ - /* 0x000fea0003800000 */ - /*00f0*/ BRA 0xf0; /* 0xfffffff000007947 */ - /* 0x000fc0000383ffff */ - ........................................... - - - Function : _Z12wmma_exampleP6__halfS0_Pfiiiff - .headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)" - /*0000*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ; /* 0x000000fffffff389 */ - /* 0x000fe200000e00ff */ - /*0010*/ IMAD.U32 R1, RZ, RZ, c[0x0][0x28]; /* 0x00000a00ff017624 */ - /* 0x000fd000078e00ff */ - /*0020*/ IADD3 R1, R1, -0x8, RZ; /* 0xfffffff801017810 */ - /* 0x000fc80007ffe0ff */ - /*0030*/ IADD3 R2, P0, R1, c[0x0][0x20], RZ; /* 0x0000080001027a10 */ - /* 0x000fca0007f1e0ff */ - /*0040*/ IMAD.X R0, RZ, RZ, c[0x0][0x24], P0; /* 0x00000900ff007624 */ - /* 0x000fd000000e06ff */ - /*0050*/ CS2R.32 R3, SR_CLOCKLO; /* 0x0000000000037805 */ - /* 0x000fd00000005000 */ - /*0060*/ I2F.U32.RP R7, 0x20; /* 0x0000002000077906 */ - /* 0x000e240000209000 */ - /*0070*/ MUFU.RCP R7, R7; /* 0x0000000700077308 */ - /* 0x001e220000001000 */ - /*0080*/ S2R R6, SR_CTAID.X; /* 0x0000000000067919 */ - /* 0x000e620000002500 */ - /*0090*/ S2R R9, SR_TID.X; /* 0x0000000000097919 */ - /* 0x000e620000002100 */ - /*00a0*/ IADD3 R8, R7, 0xffffffe, RZ; /* 0x0ffffffe07087810 */ - /* 0x001fcc0007ffe0ff */ - /*00b0*/ F2I.FTZ.U32.TRUNC.NTZ R5, R8; /* 0x0000000800057305 */ - /* 0x0000a2000021f000 */ - /*00c0*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fe400078e00ff */ - /*00d0*/ IMAD R6, R6, c[0x0][0x0], R9; /* 0x0000000006067a24 */ - /* 0x002fe400078e0209 */ - /*00e0*/ IMAD.U32 R10, R5, -0x20, RZ; /* 0xffffffe0050a7824 */ - /* 0x004fc800078e00ff */ - /*00f0*/ IMAD.WIDE.U32 R4, R5, R10, R4; /* 0x0000000a05047225 */ - /* 0x000fd000078e0004 */ - /*0100*/ IMAD.WIDE.U32 R4, R5, R6, RZ; /* 0x0000000605047225 */ - /* 0x000fcc00078e00ff */ - /*0110*/ IMAD R4, RZ, RZ, -R5; /* 0x000000ffff047224 */ - /* 0x000fc800078e0a05 */ - /*0120*/ IMAD R6, R4, 0x20, R6; /* 0x0000002004067824 */ - /* 0x000fca00078e0206 */ - /*0130*/ ISETP.GE.U32.AND P0, PT, R6.reuse, 0x20, PT, !PT; /* 0x000000200600780c */ - /* 0x040fe20003f060f0 */ - /*0140*/ S2R R28, SR_CTAID.Y; /* 0x00000000001c7919 */ - /* 0x000e220000002600 */ - /*0150*/ S2R R7, SR_TID.Y; /* 0x0000000000077919 */ - /* 0x000e340000002200 */ - /*0160*/ @P0 IADD3 R6, R6, -0x20, RZ; /* 0xffffffe006060810 */ - /* 0x000fc80007ffe0ff */ - /*0170*/ ISETP.GE.U32.AND P1, PT, R6, 0x20, PT, !PT; /* 0x000000200600780c */ - /* 0x000fe40003f260f0 */ - /*0180*/ @P0 IADD3 R5, R5, 0x1, RZ; /* 0x0000000105050810 */ - /* 0x000fe40007ffe0ff */ - /*0190*/ ISETP.LT.AND P0, PT, RZ, c[0x0][0x180], PT, !PT; /* 0x00006000ff007a0c */ - /* 0x000fd00003f012f0 */ - /*01a0*/ @P1 IADD3 R5, R5, 0x1, RZ; /* 0x0000000105051810 */ - /* 0x000fe20007ffe0ff */ - /*01b0*/ IMAD R28, R28, c[0x0][0x4], R7; /* 0x000001001c1c7a24 */ - /* 0x001fc600078e0207 */ - /*01c0*/ SHF.L.U32 R29, R5, 0x4, RZ; /* 0x00000004051d7819 */ - /* 0x000fe200000006ff */ - /*01d0*/ IMAD.U32 R28, R28, 0x10, RZ; /* 0x000000101c1c7824 */ - /* 0x000fc600078e00ff */ - /*01e0*/ ISETP.LT.AND P0, PT, R29, c[0x0][0x178], P0, !PT; /* 0x00005e001d007a0c */ - /* 0x000fc800007012f0 */ - /*01f0*/ ISETP.LT.AND P0, PT, R28, c[0x0][0x17c], P0, !PT; /* 0x00005f001c007a0c */ - /* 0x000fe200007012f0 */ - /*0200*/ BSSY B0, 0x5b0; /* 0x000003a000007945 */ - /* 0x000fe20003800000 */ - /*0210*/ IMAD.U32 R7, RZ, RZ, RZ; /* 0x000000ffff077224 */ - /* 0x000fe200078e00ff */ - /*0220*/ MOV R11, RZ; /* 0x000000ff000b7202 */ - /* 0x000fe20000000f00 */ - /*0230*/ IMAD.U32 R6, RZ, RZ, RZ; /* 0x000000ffff067224 */ - /* 0x000fe400078e00ff */ - /*0240*/ IMAD.U32 R5, RZ, RZ, RZ; /* 0x000000ffff057224 */ - /* 0x000fe400078e00ff */ - /*0250*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fe400078e00ff */ - /*0260*/ IMAD.U32 R10, RZ, RZ, RZ; /* 0x000000ffff0a7224 */ - /* 0x000fc400078e00ff */ - /*0270*/ IMAD.U32 R9, RZ, RZ, RZ; /* 0x000000ffff097224 */ - /* 0x000fe400078e00ff */ - /*0280*/ IMAD.U32 R8, RZ, RZ, RZ; /* 0x000000ffff087224 */ - /* 0x000fe200078e00ff */ - /*0290*/ @!P0 BRA 0x5a0; /* 0x0000030000008947 */ - /* 0x000fee0003800000 */ - /*02a0*/ S2R R6, SR_LANEID; /* 0x0000000000067919 */ - /* 0x000e220000000000 */ - /*02b0*/ IMAD.U32 R10, RZ, RZ, 0x2; /* 0x00000002ff0a7424 */ - /* 0x000fc800078e00ff */ - /*02c0*/ IMAD.WIDE R16, R29, R10, c[0x0][0x160]; /* 0x000058001d107625 */ - /* 0x000fe200078e020a */ - /*02d0*/ SHF.R.U32.HI R4, RZ, 0x2, R6; /* 0x00000002ff047819 */ - /* 0x001fc80000011606 */ - /*02e0*/ LOP3.LUT R5, R4, 0x3, RZ, 0xc0, !PT; /* 0x0000000304057812 */ - /* 0x000fe400078ec0ff */ - /*02f0*/ LOP3.LUT R4, R6, 0x3, RZ, 0xc0, !PT; /* 0x0000000306047812 */ - /* 0x000fe400078ec0ff */ - /*0300*/ LOP3.LUT R7, R5, 0x1, RZ, 0xc0, !PT; /* 0x0000000105077812 */ - /* 0x000fe400078ec0ff */ - /*0310*/ SHF.R.U32.HI R6, RZ, 0x4, R6; /* 0x00000004ff067819 */ - /* 0x000fe40000011606 */ - /*0320*/ SHF.R.U32.HI R5, RZ, 0x1, R5; /* 0x00000001ff057819 */ - /* 0x000fe20000011605 */ - /*0330*/ IMAD R7, R7, 0x8, R4; /* 0x0000000807077824 */ - /* 0x000fe200078e0204 */ - /*0340*/ LOP3.LUT R6, R6, 0x1, RZ, 0xc0, !PT; /* 0x0000000106067812 */ - /* 0x000fc400078ec0ff */ - /*0350*/ LEA R4, R5, R4, 0x3; /* 0x0000000405047211 */ - /* 0x000fe200078e18ff */ - /*0360*/ IMAD.WIDE R12, R28, R10, c[0x0][0x168]; /* 0x00005a001c0c7625 */ - /* 0x000fe400078e020a */ - /*0370*/ IMAD R7, R6.reuse, 0x4, R7; /* 0x0000000406077824 */ - /* 0x040fe400078e0207 */ - /*0380*/ IMAD R4, R6, 0x4, R4; /* 0x0000000406047824 */ - /* 0x000fe400078e0204 */ - /*0390*/ IMAD.U32 R7, R7, 0x2, RZ; /* 0x0000000207077824 */ - /* 0x000fe400078e00ff */ - /*03a0*/ IMAD.U32 R5, R4, 0x2, RZ; /* 0x0000000204057824 */ - /* 0x000fe400078e00ff */ - /*03b0*/ IMAD.WIDE.U32 R16, R7, c[0x0][0x178], R16; /* 0x00005e0007107a25 */ - /* 0x000fc400078e0010 */ - /*03c0*/ IMAD.WIDE.U32 R12, R5, c[0x0][0x180], R12; /* 0x00006000050c7a25 */ - /* 0x000fd000078e000c */ - /*03d0*/ LD.E.128.SYS R24, [R16]; /* 0x0000000010187980 */ - /* 0x000064000010ed00 */ - /*03e0*/ LD.E.128.SYS R20, [R12]; /* 0x000000000c147980 */ - /* 0x000462000010ed00 */ - /*03f0*/ LD.E.128.SYS R16, [R16+0x10]; /* 0x0000001010107980 */ - /* 0x001e22000010ed00 */ - /*0400*/ LD.E.128.SYS R12, [R12+0x10]; /* 0x000000100c0c7980 */ - /* 0x004e22000010ed00 */ - /*0410*/ IMAD.U32 R8, RZ, RZ, RZ; /* 0x000000ffff087224 */ - /* 0x000fe200078e00ff */ - /*0420*/ MOV R9, RZ; /* 0x000000ff00097202 */ - /* 0x000fe20000000f00 */ - /*0430*/ IMAD.U32 R10, RZ, RZ, RZ; /* 0x000000ffff0a7224 */ - /* 0x000fe400078e00ff */ - /*0440*/ IMAD.U32 R11, RZ, RZ, RZ; /* 0x000000ffff0b7224 */ - /* 0x000fe200078e00ff */ - /*0450*/ MOV R7, RZ; /* 0x000000ff00077202 */ - /* 0x000fe20000000f00 */ - /*0460*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fc400078e00ff */ - /*0470*/ IMAD.U32 R5, RZ, RZ, RZ; /* 0x000000ffff057224 */ - /* 0x000fe400078e00ff */ - /*0480*/ IMAD.U32 R6, RZ, RZ, RZ; /* 0x000000ffff067224 */ - /* 0x000fe200078e00ff */ - /*0490*/ WARPSYNC 0xffffffff; /* 0xffffffff00007948 */ - /* 0x000fe20003800000 */ - /*04a0*/ HMMA.884.F32.F32.STEP0 R8, R24.reuse, R20.reuse.T, R8; /* 0x0000001418087236 */ - /* 0x0c22640000005408 */ - /*04b0*/ HMMA.884.F32.F32.STEP1 R10, R24.reuse, R20.reuse.T, R10; /* 0x00000014180a7236 */ - /* 0x0c04a4000000d40a */ - /*04c0*/ HMMA.884.F32.F32.STEP2 R4, R24.reuse, R20.reuse.T, R4; /* 0x0000001418047236 */ - /* 0x0c06e40000015404 */ - /*04d0*/ HMMA.884.F32.F32.STEP3 R6, R24, R20.T, R6; /* 0x0000001418067236 */ - /* 0x000928000001d406 */ - /*04e0*/ HMMA.884.F32.F32.STEP0 R8, R26.reuse, R22.reuse.T, R8; /* 0x000000161a087236 */ - /* 0x0c20240000005408 */ - /*04f0*/ HMMA.884.F32.F32.STEP1 R10, R26.reuse, R22.reuse.T, R10; /* 0x000000161a0a7236 */ - /* 0x0c4264000000d40a */ - /*0500*/ HMMA.884.F32.F32.STEP2 R4, R26.reuse, R22.reuse.T, R4; /* 0x000000161a047236 */ - /* 0x0c84a40000015404 */ - /*0510*/ HMMA.884.F32.F32.STEP3 R6, R26, R22.T, R6; /* 0x000000161a067236 */ - /* 0x0106e8000001d406 */ - /*0520*/ HMMA.884.F32.F32.STEP0 R8, R16.reuse, R12.reuse.T, R8; /* 0x0000000c10087236 */ - /* 0x0c10240000005408 */ - /*0530*/ HMMA.884.F32.F32.STEP1 R10, R16.reuse, R12.reuse.T, R10; /* 0x0000000c100a7236 */ - /* 0x0c2264000000d40a */ - /*0540*/ HMMA.884.F32.F32.STEP2 R4, R16.reuse, R12.reuse.T, R4; /* 0x0000000c10047236 */ - /* 0x0c44a40000015404 */ - /*0550*/ HMMA.884.F32.F32.STEP3 R6, R16, R12.T, R6; /* 0x0000000c10067236 */ - /* 0x0086e8000001d406 */ - /*0560*/ HMMA.884.F32.F32.STEP0 R8, R18.reuse, R14.reuse.T, R8; /* 0x0000000e12087236 */ - /* 0x0c10240000005408 */ - /*0570*/ HMMA.884.F32.F32.STEP1 R10, R18.reuse, R14.reuse.T, R10; /* 0x0000000e120a7236 */ - /* 0x0c2024000000d40a */ - /*0580*/ HMMA.884.F32.F32.STEP2 R4, R18.reuse, R14.reuse.T, R4; /* 0x0000000e12047236 */ - /* 0x0c40240000015404 */ - /*0590*/ HMMA.884.F32.F32.STEP3 R6, R18, R14.T, R6; /* 0x0000000e12067236 */ - /* 0x008034000001d406 */ - /*05a0*/ BSYNC B0; /* 0x0000000000007941 */ - /* 0x001fea0003800000 */ - /*05b0*/ S2R R12, SR_LANEID; /* 0x00000000000c7919 */ - /* 0x000e220000000000 */ - /*05c0*/ IMAD R28, R28, c[0x0][0x178], RZ; /* 0x00005e001c1c7a24 */ - /* 0x000fe200078e02ff */ - /*05d0*/ SHF.R.U32.HI R14, RZ, 0x4, R12; /* 0x00000004ff0e7819 */ - /* 0x001fc4000001160c */ - /*05e0*/ SHF.R.U32.HI R13, RZ, 0x2, R12; /* 0x00000002ff0d7819 */ - /* 0x000fe4000001160c */ - /*05f0*/ LOP3.LUT R12, R12, 0x3, RZ, 0xc0, !PT; /* 0x000000030c0c7812 */ - /* 0x000fe400078ec0ff */ - /*0600*/ LOP3.LUT R14, R14, 0x1, RZ, 0xc0, !PT; /* 0x000000010e0e7812 */ - /* 0x000fe400078ec0ff */ - /*0610*/ LOP3.LUT R13, R13, 0x3, RZ, 0xc0, !PT; /* 0x000000030d0d7812 */ - /* 0x000fc600078ec0ff */ - /*0620*/ IMAD R12, R14, 0x4, R12; /* 0x000000040e0c7824 */ - /* 0x000fe200078e020c */ - /*0630*/ LOP3.LUT R15, R13, 0x1, RZ, 0xc0, !PT; /* 0x000000010d0f7812 */ - /* 0x000fe400078ec0ff */ - /*0640*/ SHF.R.U32.HI R16, RZ, 0x1, R13; /* 0x00000001ff107819 */ - /* 0x000fe4000001160d */ - /*0650*/ LOP3.LUT R13, R12.reuse, 0x5, RZ, 0xc0, !PT; /* 0x000000050c0d7812 */ - /* 0x040fe400078ec0ff */ - /*0660*/ LOP3.LUT R14, R12, 0x2, RZ, 0xc0, !PT; /* 0x000000020c0e7812 */ - /* 0x000fc600078ec0ff */ - /*0670*/ IMAD R12, R15, 0x8, R13; /* 0x000000080f0c7824 */ - /* 0x000fe200078e020d */ - /*0680*/ SHF.R.S32.HI R15, RZ, 0x1f, R29; /* 0x0000001fff0f7819 */ - /* 0x000fe2000001141d */ - /*0690*/ IMAD R14, R16, 0x8, R14; /* 0x00000008100e7824 */ - /* 0x000fe200078e020e */ - /*06a0*/ IADD3 R17, P0, R29, R28, RZ; /* 0x0000001c1d117210 */ - /* 0x000fe20007f1e0ff */ - /*06b0*/ IMAD.U32 R13, RZ, RZ, RZ; /* 0x000000ffff0d7224 */ - /* 0x000fc600078e00ff */ - /*06c0*/ LEA.HI.X.SX32 R15, R28, R15, 0x1, P0; /* 0x0000000f1c0f7211 */ - /* 0x000fe200000f0eff */ - /*06d0*/ IMAD.WIDE.U32 R12, R14, c[0x0][0x178], R12; /* 0x00005e000e0c7a25 */ - /* 0x000fe200078e000c */ - /*06e0*/ LEA R14, P0, R17, c[0x0][0x170], 0x2; /* 0x00005c00110e7a11 */ - /* 0x000fe200078010ff */ - /*06f0*/ IMAD.U32 R18, RZ, RZ, c[0x0][0x178]; /* 0x00005e00ff127624 */ - /* 0x000fc600078e00ff */ - /*0700*/ LEA.HI.X R17, R17, c[0x0][0x174], R15, 0x2, P0; /* 0x00005d0011117a11 */ - /* 0x000fe400000f140f */ - /*0710*/ LEA R16, P0, R12, R14, 0x2; /* 0x0000000e0c107211 */ - /* 0x000fe400078010ff */ - /*0720*/ SHF.L.U32 R14, R18, 0x2, RZ; /* 0x00000002120e7819 */ - /* 0x000fe400000006ff */ - /*0730*/ SHF.R.U32.HI R15, RZ, 0x1e, R18; /* 0x0000001eff0f7819 */ - /* 0x000fe40000011612 */ - /*0740*/ LEA.HI.X R13, R12, R17, R13, 0x2, P0; /* 0x000000110c0d7211 */ - /* 0x000fe400000f140d */ - /*0750*/ LEA R18, P1, R14, R16, 0x2; /* 0x000000100e127211 */ - /* 0x000fc400078210ff */ - /*0760*/ IADD3 R17, P0, R14.reuse, R16, RZ; /* 0x000000100e117210 */ - /* 0x040fe40007f1e0ff */ - /*0770*/ LEA.HI.X R19, R14.reuse, R13, R15, 0x2, P1; /* 0x0000000d0e137211 */ - /* 0x040fe400008f140f */ - /*0780*/ IADD3 R21, P1, R14, R18, RZ; /* 0x000000120e157210 */ - /* 0x000fe20007f3e0ff */ - /*0790*/ IMAD.X R20, R15.reuse, 0x1, R13, P0; /* 0x000000010f147824 */ - /* 0x040fe400000e060d */ - /*07a0*/ IMAD.U32 R12, RZ, RZ, R16; /* 0x000000ffff0c7224 */ - /* 0x000fe400078e0010 */ - /*07b0*/ IMAD.X R22, R15, 0x1, R19, P1; /* 0x000000010f167824 */ - /* 0x000fe400008e0613 */ - /*07c0*/ IMAD.U32 R14, RZ, RZ, R17; /* 0x000000ffff0e7224 */ - /* 0x000fc400078e0011 */ - /*07d0*/ IMAD.U32 R15, RZ, RZ, R20; /* 0x000000ffff0f7224 */ - /* 0x000fe200078e0014 */ - /*07e0*/ MOV R16, R21; /* 0x0000001500107202 */ - /* 0x000fe20000000f00 */ - /*07f0*/ IMAD.U32 R17, RZ, RZ, R22; /* 0x000000ffff117224 */ - /* 0x000fe200078e0016 */ - /*0800*/ ST.E.SYS [R12], R8; /* 0x000000000c007385 */ - /* 0x0001e2000010e908 */ - /*0810*/ ST.E.SYS [R12+0x8], R10; /* 0x000000080c007385 */ - /* 0x0003e8000010e90a */ - /*0820*/ ST.E.SYS [R14], R9; /* 0x000000000e007385 */ - /* 0x0003e2000010e909 */ - /*0830*/ ST.E.SYS [R14+0x8], R11; /* 0x000000080e007385 */ - /* 0x0003e2000010e90b */ - /*0840*/ ST.E.SYS [R18], R4; /* 0x0000000012007385 */ - /* 0x0003e2000010e904 */ - /*0850*/ ST.E.SYS [R18+0x8], R6; /* 0x0000000812007385 */ - /* 0x0003e2000010e906 */ - /*0860*/ ST.E.SYS [R16], R5; /* 0x0000000010007385 */ - /* 0x0003e2000010e905 */ - /*0870*/ ST.E.SYS [R16+0x8], R7; /* 0x0000000810007385 */ - /* 0x0003e2000010e907 */ - /*0880*/ WARPSYNC 0xffffffff; /* 0xffffffff00007948 */ - /* 0x000fe20003800000 */ - /*0890*/ IADD3 R8, R2, -c[0x0][0x20], RZ; /* 0x8000080002087a10 */ - /* 0x001fd00007ffe0ff */ - /*08a0*/ CS2R.32 R4, SR_CLOCKLO; /* 0x0000000000047805 */ - /* 0x002fd00000005000 */ - /*08b0*/ IMAD R3, R4, 0x1, -R3; /* 0x0000000104037824 */ - /* 0x000fd000078e0a03 */ - /*08c0*/ STL [R8], R3; /* 0x0000000308007387 */ - /* 0x0001e20000100800 */ - /*08d0*/ IMAD.U32 R6, RZ, RZ, R2; /* 0x000000ffff067224 */ - /* 0x000fe200078e0002 */ - /*08e0*/ MOV R4, 0x0; /* 0x0000000000047802 */ - /* 0x000fe20000000f00 */ - /*08f0*/ IMAD.U32 R7, RZ, RZ, R0; /* 0x000000ffff077224 */ - /* 0x000fe200078e0000 */ - /*0900*/ MOV R5, 0x0; /* 0x0000000000057802 */ - /* 0x000fe40000000f00 */ - /*0910*/ MOV R20, 0x0; /* 0x0000000000147802 */ - /* 0x000fe40000000f00 */ - /*0920*/ MOV R21, 0x0; /* 0x0000000000157802 */ - /* 0x000fd00000000f00 */ - /*0930*/ CALL.ABS.NOINC 0x0; /* 0x0000000000007943 */ - /* 0x001fea0003c00000 */ - /*0940*/ EXIT; /* 0x000000000000794d */ - /* 0x000fea0003800000 */ - /*0950*/ BRA 0x950; /* 0xfffffff000007947 */ - /* 0x000fc0000383ffff */ - /*0960*/ NOP; /* 0x0000000000007918 */ - /* 0x000fc00000000000 */ - /*0970*/ NOP; /* 0x0000000000007918 */ - /* 0x000fc00000000000 */ - ............................................. - - - diff --git a/cuda-kernels/_cuobjdump_complete_output_EIGzTK b/cuda-kernels/_cuobjdump_complete_output_EIGzTK deleted file mode 100644 index 36999c0..0000000 --- a/cuda-kernels/_cuobjdump_complete_output_EIGzTK +++ /dev/null @@ -1,1055 +0,0 @@ - -Fatbin elf code: -================ -arch = sm_70 -code version = [1,7] -producer = -host = linux -compile_size = 64bit - -64bit elf: type=2, abi=7, sm=70, toolkit=90, flags = 0x460546 -Sections: -Index Offset Size ES Align Type Flags Link Info Name - 1 40 32 0 1 STRTAB 0 0 0 .shstrtab - 2 72 32 0 1 STRTAB 0 0 0 .strtab - 3 a8 18 18 8 SYMTAB 0 2 0 .symtab - -.section .strtab - -.section .shstrtab - -.section .symtab - index value size info other shndx name - 0 0 0 0 0 0 (null) - - code for sm_70 - -Fatbin elf code: -================ -arch = sm_70 -code version = [1,7] -producer = cuda -host = linux -compile_size = 64bit - -64bit elf: type=2, abi=7, sm=70, toolkit=90, flags = 0x460546 -Sections: -Index Offset Size ES Align Type Flags Link Info Name - 1 40 21b 0 1 STRTAB 0 0 0 .shstrtab - 2 25b 273 0 1 STRTAB 0 0 0 .strtab - 3 4d0 108 18 8 SYMTAB 0 2 7 .symtab - 4 5d8 e0 0 1 PROGBITS 0 0 0 .debug_frame - 5 6b8 48 0 4 CUDA_INFO 0 3 0 .nv.info - 6 700 50 0 4 CUDA_INFO 0 3 d .nv.info._Z17convertFp32ToFp16P6__halfPfi - 7 750 ac 0 4 CUDA_INFO 0 3 e .nv.info._Z12wmma_exampleP6__halfS0_Pfiiiff - 8 800 30 10 8 REL 0 3 e .rel.text._Z12wmma_exampleP6__halfS0_Pfiiiff - 9 830 30 18 8 RELA 0 3 e .rela.text._Z12wmma_exampleP6__halfS0_Pfiiiff - a 860 20 10 8 REL 0 3 4 .rel.debug_frame - b 880 174 0 4 PROGBITS 2 0 d .nv.constant0._Z17convertFp32ToFp16P6__halfPfi - c 9f4 18c 0 4 PROGBITS 2 0 e .nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff - d b80 100 0 80 PROGBITS 6 3 9000008 .text._Z17convertFp32ToFp16P6__halfPfi - e c80 980 0 80 PROGBITS 6 3 20000009 .text._Z12wmma_exampleP6__halfS0_Pfiiiff - f 1600 9 0 10 PROGBITS 3 0 0 .nv.global.init - -.section .strtab - -.section .shstrtab - -.section .symtab - index value size info other shndx name - 0 0 0 0 0 0 (null) - 1 0 0 3 0 d .text._Z17convertFp32ToFp16P6__halfPfi - 2 0 0 3 0 f .nv.global.init - 3 0 9 1 0 f $str - 4 0 0 3 0 b .nv.constant0._Z17convertFp32ToFp16P6__halfPfi - 5 0 0 3 0 e .text._Z12wmma_exampleP6__halfS0_Pfiiiff - 6 0 0 3 0 c .nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff - 7 0 0 3 0 4 .debug_frame - 8 0 256 12 10 d _Z17convertFp32ToFp16P6__halfPfi - 9 0 2432 12 10 e _Z12wmma_exampleP6__halfS0_Pfiiiff - 10 0 0 12 0 0 vprintf - - -.nv.constant0._Z17convertFp32ToFp16P6__halfPfi -0x00000000 0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 - - - -.nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff -0x00000000 0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 - - -.nv.global.init -0x636f6c63 0x64253d6b 0 - - -.nv.info - <0x1> - Attribute: EIATTR_MAX_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x9 0x0 - <0x2> - Attribute: EIATTR_MIN_STACK_SIZE - Format: EIFMT_SVAL - Value: function: _Z12wmma_exampleP6__halfS0_Pfiiiff(0x9) min stack size: 0x8 - <0x3> - Attribute: EIATTR_FRAME_SIZE - Format: EIFMT_SVAL - Value: function: _Z12wmma_exampleP6__halfS0_Pfiiiff(0x9) frame size: 0x8 - <0x4> - Attribute: EIATTR_MAX_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x8 0x0 - <0x5> - Attribute: EIATTR_MIN_STACK_SIZE - Format: EIFMT_SVAL - Value: function: _Z17convertFp32ToFp16P6__halfPfi(0x8) min stack size: 0x0 - <0x6> - Attribute: EIATTR_FRAME_SIZE - Format: EIFMT_SVAL - Value: function: _Z17convertFp32ToFp16P6__halfPfi(0x8) frame size: 0x0 - - -.nv.info._Z17convertFp32ToFp16P6__halfPfi - <0x1> - Attribute: EIATTR_PARAM_CBANK - Format: EIFMT_SVAL - Value: 0x4 0x140160 - <0x2> - Attribute: EIATTR_CBANK_PARAM_SIZE - Format: EIFMT_HVAL - Value: 0x14 - <0x3> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x2 Offset : 0x10 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x4> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x1 Offset : 0x8 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x5> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x0 Offset : 0x0 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x6> - Attribute: EIATTR_MAXREG_COUNT - Format: EIFMT_HVAL - Value: 0xff - <0x7> - Attribute: EIATTR_EXIT_INSTR_OFFSETS - Format: EIFMT_SVAL - Value: 0x60 0xe0 - - -.nv.info._Z12wmma_exampleP6__halfS0_Pfiiiff - <0x1> - Attribute: EIATTR_PARAM_CBANK - Format: EIFMT_SVAL - Value: 0x6 0x2c0160 - <0x2> - Attribute: EIATTR_CBANK_PARAM_SIZE - Format: EIFMT_HVAL - Value: 0x2c - <0x3> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x7 Offset : 0x28 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x4> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x6 Offset : 0x24 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x5> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x5 Offset : 0x20 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x6> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x4 Offset : 0x1c Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x7> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x3 Offset : 0x18 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x8> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x2 Offset : 0x10 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x9> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x1 Offset : 0x8 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x10> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x0 Offset : 0x0 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x11> - Attribute: EIATTR_MAXREG_COUNT - Format: EIFMT_HVAL - Value: 0xff - <0x12> - Attribute: EIATTR_EXIT_INSTR_OFFSETS - Format: EIFMT_SVAL - Value: 0x940 - <0x13> - Attribute: EIATTR_EXTERNS - Format: EIFMT_SVAL - Value: externs: vprintf(0xa) - <0x14> - Attribute: EIATTR_CRS_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x0 - - -.text._Z17convertFp32ToFp16P6__halfPfi -bar = 0 reg = 9 lmem=0 smem=0 -0xfffff389 0x000000ff 0x000e00ff 0x000fe200 -0x00017a02 0x00000a00 0x00000f00 0x000fd000 -0x00047919 0x00000000 0x00002500 0x000e2200 -0x00027919 0x00000000 0x00002100 0x000e2400 -0x04047a24 0x00000000 0x078e0202 0x001fca00 -0x04007a0c 0x00005c00 0x03f062f0 0x000fd800 -0x0000094d 0x00000000 0x03800000 0x000fea00 -0x00027802 0x00000004 0x00000f00 0x000fca00 -0x04027625 0x00005a00 0x078e0202 0x000fd400 -0x02027381 0x00000000 0x001ee900 0x000e2200 -0x00057802 0x00000002 0x00000f00 0x000fca00 -0x04047625 0x00005800 0x078e0205 0x000fe200 -0x00067304 0x00000002 0x00200800 0x001e3200 -0x04007386 0x00000006 0x0010e500 0x0011e200 -0x0000794d 0x00000000 0x03800000 0x000fea00 -0x00007947 0xfffffff0 0x0383ffff 0x000fc000 - - - -.text._Z12wmma_exampleP6__halfS0_Pfiiiff -bar = 0 reg = 32 lmem=0 smem=0 -0xfffff389 0x000000ff 0x000e00ff 0x000fe200 -0xff017624 0x00000a00 0x078e00ff 0x000fd000 -0x01017810 0xfffffff8 0x07ffe0ff 0x000fc800 -0x01027a10 0x00000800 0x07f1e0ff 0x000fca00 -0xff007624 0x00000900 0x000e06ff 0x000fd000 -0x00037805 0x00000000 0x00005000 0x000fd000 -0x00077906 0x00000020 0x00209000 0x000e2400 -0x00077308 0x00000007 0x00001000 0x001e2200 -0x00067919 0x00000000 0x00002500 0x000e6200 -0x00097919 0x00000000 0x00002100 0x000e6200 -0x07087810 0x0ffffffe 0x07ffe0ff 0x001fcc00 -0x00057305 0x00000008 0x0021f000 0x0000a200 -0xff047224 0x000000ff 0x078e00ff 0x000fe400 -0x06067a24 0x00000000 0x078e0209 0x002fe400 -0x050a7824 0xffffffe0 0x078e00ff 0x004fc800 -0x05047225 0x0000000a 0x078e0004 0x000fd000 -0x05047225 0x00000006 0x078e00ff 0x000fcc00 -0xff047224 0x000000ff 0x078e0a05 0x000fc800 -0x04067824 0x00000020 0x078e0206 0x000fca00 -0x0600780c 0x00000020 0x03f060f0 0x040fe200 -0x001c7919 0x00000000 0x00002600 0x000e2200 -0x00077919 0x00000000 0x00002200 0x000e3400 -0x06060810 0xffffffe0 0x07ffe0ff 0x000fc800 -0x0600780c 0x00000020 0x03f260f0 0x000fe400 -0x05050810 0x00000001 0x07ffe0ff 0x000fe400 -0xff007a0c 0x00006000 0x03f012f0 0x000fd000 -0x05051810 0x00000001 0x07ffe0ff 0x000fe200 -0x1c1c7a24 0x00000100 0x078e0207 0x001fc600 -0x051d7819 0x00000004 0x000006ff 0x000fe200 -0x1c1c7824 0x00000010 0x078e00ff 0x000fc600 -0x1d007a0c 0x00005e00 0x007012f0 0x000fc800 -0x1c007a0c 0x00005f00 0x007012f0 0x000fe200 -0x00007945 0x000003a0 0x03800000 0x000fe200 -0xff077224 0x000000ff 0x078e00ff 0x000fe200 -0x000b7202 0x000000ff 0x00000f00 0x000fe200 -0xff067224 0x000000ff 0x078e00ff 0x000fe400 -0xff057224 0x000000ff 0x078e00ff 0x000fe400 -0xff047224 0x000000ff 0x078e00ff 0x000fe400 -0xff0a7224 0x000000ff 0x078e00ff 0x000fc400 -0xff097224 0x000000ff 0x078e00ff 0x000fe400 -0xff087224 0x000000ff 0x078e00ff 0x000fe200 -0x00008947 0x00000300 0x03800000 0x000fee00 -0x00067919 0x00000000 0x00000000 0x000e2200 -0xff0a7424 0x00000002 0x078e00ff 0x000fc800 -0x1d107625 0x00005800 0x078e020a 0x000fe200 -0xff047819 0x00000002 0x00011606 0x001fc800 -0x04057812 0x00000003 0x078ec0ff 0x000fe400 -0x06047812 0x00000003 0x078ec0ff 0x000fe400 -0x05077812 0x00000001 0x078ec0ff 0x000fe400 -0xff067819 0x00000004 0x00011606 0x000fe400 -0xff057819 0x00000001 0x00011605 0x000fe200 -0x07077824 0x00000008 0x078e0204 0x000fe200 -0x06067812 0x00000001 0x078ec0ff 0x000fc400 -0x05047211 0x00000004 0x078e18ff 0x000fe200 -0x1c0c7625 0x00005a00 0x078e020a 0x000fe400 -0x06077824 0x00000004 0x078e0207 0x040fe400 -0x06047824 0x00000004 0x078e0204 0x000fe400 -0x07077824 0x00000002 0x078e00ff 0x000fe400 -0x04057824 0x00000002 0x078e00ff 0x000fe400 -0x07107a25 0x00005e00 0x078e0010 0x000fc400 -0x050c7a25 0x00006000 0x078e000c 0x000fd000 -0x10187980 0x00000000 0x0010ed00 0x00006400 -0x0c147980 0x00000000 0x0010ed00 0x00046200 -0x10107980 0x00000010 0x0010ed00 0x001e2200 -0x0c0c7980 0x00000010 0x0010ed00 0x004e2200 -0xff087224 0x000000ff 0x078e00ff 0x000fe200 -0x00097202 0x000000ff 0x00000f00 0x000fe200 -0xff0a7224 0x000000ff 0x078e00ff 0x000fe400 -0xff0b7224 0x000000ff 0x078e00ff 0x000fe200 -0x00077202 0x000000ff 0x00000f00 0x000fe200 -0xff047224 0x000000ff 0x078e00ff 0x000fc400 -0xff057224 0x000000ff 0x078e00ff 0x000fe400 -0xff067224 0x000000ff 0x078e00ff 0x000fe200 -0x00007948 0xffffffff 0x03800000 0x000fe200 -0x18087236 0x00000014 0x00005408 0x0c226400 -0x180a7236 0x00000014 0x0000d40a 0x0c04a400 -0x18047236 0x00000014 0x00015404 0x0c06e400 -0x18067236 0x00000014 0x0001d406 0x00092800 -0x1a087236 0x00000016 0x00005408 0x0c202400 -0x1a0a7236 0x00000016 0x0000d40a 0x0c426400 -0x1a047236 0x00000016 0x00015404 0x0c84a400 -0x1a067236 0x00000016 0x0001d406 0x0106e800 -0x10087236 0x0000000c 0x00005408 0x0c102400 -0x100a7236 0x0000000c 0x0000d40a 0x0c226400 -0x10047236 0x0000000c 0x00015404 0x0c44a400 -0x10067236 0x0000000c 0x0001d406 0x0086e800 -0x12087236 0x0000000e 0x00005408 0x0c102400 -0x120a7236 0x0000000e 0x0000d40a 0x0c202400 -0x12047236 0x0000000e 0x00015404 0x0c402400 -0x12067236 0x0000000e 0x0001d406 0x00803400 -0x00007941 0x00000000 0x03800000 0x001fea00 -0x000c7919 0x00000000 0x00000000 0x000e2200 -0x1c1c7a24 0x00005e00 0x078e02ff 0x000fe200 -0xff0e7819 0x00000004 0x0001160c 0x001fc400 -0xff0d7819 0x00000002 0x0001160c 0x000fe400 -0x0c0c7812 0x00000003 0x078ec0ff 0x000fe400 -0x0e0e7812 0x00000001 0x078ec0ff 0x000fe400 -0x0d0d7812 0x00000003 0x078ec0ff 0x000fc600 -0x0e0c7824 0x00000004 0x078e020c 0x000fe200 -0x0d0f7812 0x00000001 0x078ec0ff 0x000fe400 -0xff107819 0x00000001 0x0001160d 0x000fe400 -0x0c0d7812 0x00000005 0x078ec0ff 0x040fe400 -0x0c0e7812 0x00000002 0x078ec0ff 0x000fc600 -0x0f0c7824 0x00000008 0x078e020d 0x000fe200 -0xff0f7819 0x0000001f 0x0001141d 0x000fe200 -0x100e7824 0x00000008 0x078e020e 0x000fe200 -0x1d117210 0x0000001c 0x07f1e0ff 0x000fe200 -0xff0d7224 0x000000ff 0x078e00ff 0x000fc600 -0x1c0f7211 0x0000000f 0x000f0eff 0x000fe200 -0x0e0c7a25 0x00005e00 0x078e000c 0x000fe200 -0x110e7a11 0x00005c00 0x078010ff 0x000fe200 -0xff127624 0x00005e00 0x078e00ff 0x000fc600 -0x11117a11 0x00005d00 0x000f140f 0x000fe400 -0x0c107211 0x0000000e 0x078010ff 0x000fe400 -0x120e7819 0x00000002 0x000006ff 0x000fe400 -0xff0f7819 0x0000001e 0x00011612 0x000fe400 -0x0c0d7211 0x00000011 0x000f140d 0x000fe400 -0x0e127211 0x00000010 0x078210ff 0x000fc400 -0x0e117210 0x00000010 0x07f1e0ff 0x040fe400 -0x0e137211 0x0000000d 0x008f140f 0x040fe400 -0x0e157210 0x00000012 0x07f3e0ff 0x000fe200 -0x0f147824 0x00000001 0x000e060d 0x040fe400 -0xff0c7224 0x000000ff 0x078e0010 0x000fe400 -0x0f167824 0x00000001 0x008e0613 0x000fe400 -0xff0e7224 0x000000ff 0x078e0011 0x000fc400 -0xff0f7224 0x000000ff 0x078e0014 0x000fe200 -0x00107202 0x00000015 0x00000f00 0x000fe200 -0xff117224 0x000000ff 0x078e0016 0x000fe200 -0x0c007385 0x00000000 0x0010e908 0x0001e200 -0x0c007385 0x00000008 0x0010e90a 0x0003e800 -0x0e007385 0x00000000 0x0010e909 0x0003e200 -0x0e007385 0x00000008 0x0010e90b 0x0003e200 -0x12007385 0x00000000 0x0010e904 0x0003e200 -0x12007385 0x00000008 0x0010e906 0x0003e200 -0x10007385 0x00000000 0x0010e905 0x0003e200 -0x10007385 0x00000008 0x0010e907 0x0003e200 -0x00007948 0xffffffff 0x03800000 0x000fe200 -0x02087a10 0x80000800 0x07ffe0ff 0x001fd000 -0x00047805 0x00000000 0x00005000 0x002fd000 -0x04037824 0x00000001 0x078e0a03 0x000fd000 -0x08007387 0x00000003 0x00100800 0x0001e200 -0xff067224 0x000000ff 0x078e0002 0x000fe200 -0x00047802 0x00000000 0x00000f00 0x000fe200 -0xff077224 0x000000ff 0x078e0000 0x000fe200 -0x00057802 0x00000000 0x00000f00 0x000fe400 -0x00147802 0x00000000 0x00000f00 0x000fe400 -0x00157802 0x00000000 0x00000f00 0x000fd000 -0x00007943 0x00000000 0x03c00000 0x001fea00 -0x0000794d 0x00000000 0x03800000 0x000fea00 -0x00007947 0xfffffff0 0x0383ffff 0x000fc000 -0x00007918 0x00000000 0x00000000 0x000fc000 -0x00007918 0x00000000 0x00000000 0x000fc000 - - -.section .rel.text._Z12wmma_exampleP6__halfS0_Pfiiiff REL -2272 $str R_CUDA_ABS32_LO_32 -2304 $str R_CUDA_ABS32_HI_32 -2352 vprintf R_CUDA_ABS47_34 - -.section .rela.text._Z12wmma_exampleP6__halfS0_Pfiiiff RELA -2320 _Z12wmma_exampleP6__halfS0_Pfiiiff R_CUDA_ABS32_LO_32 2368 -2336 _Z12wmma_exampleP6__halfS0_Pfiiiff R_CUDA_ABS32_HI_32 2368 - -.section .debug_frame -decodeDebugFrame, frameBuf 0xffffffff, total_length 224 -CIE length 40, cie_id -1 -version 3 -augmentation slen 1 -augmentation -code_align_factor slen 1 -data_align_factor slen 1 - Debug Frame Common Information Entry - length: 40 - CIE_id : -1 - version: 3 - augmentation: - code align factor: 4 - data align factor: -4 - return address register 16777215 - initial instructions: 23 bytes, ptr = 0x8080810c, frameBuf = 0xffffffff - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_same_value R255 - DW_CFA_same_value R1 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - Debug Frame Description Entry - length: 48 - CIE_pointer: 0 - initial_location: 0x0 - address_range: 0x100 - instructions: 24 bytes - DW_CFA_advance_loc4 delta 4 - DW_CFA_advance_loc4 delta 0 - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_advance_loc4 delta 52 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop -CIE length 40, cie_id -1 -version 3 -augmentation slen 1 -augmentation -code_align_factor slen 1 -data_align_factor slen 1 - Debug Frame Common Information Entry - length: 40 - CIE_id : -1 - version: 3 - augmentation: - code align factor: 4 - data align factor: -4 - return address register 16777215 - initial instructions: 23 bytes, ptr = 0x8080810c, frameBuf = 0xffffffff - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_same_value R255 - DW_CFA_same_value R1 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - Debug Frame Description Entry - length: 48 - CIE_pointer: 0 - initial_location: 0x0 - address_range: 0x970 - instructions: 24 bytes - DW_CFA_advance_loc4 delta 4 - DW_CFA_advance_loc4 delta 2 - DW_CFA_def_cfa register R1, offset 8 - DW_CFA_advance_loc4 delta 586 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - -.section .rel.debug_frame REL -72 _Z17convertFp32ToFp16P6__halfPfi R_NV_64 -184 _Z12wmma_exampleP6__halfS0_Pfiiiff R_NV_64 - - code for sm_70 - Function : _Z17convertFp32ToFp16P6__halfPfi - .headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)" - /*0000*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ; /* 0x000000fffffff389 */ - /* 0x000fe200000e00ff */ - /*0010*/ MOV R1, c[0x0][0x28]; /* 0x00000a0000017a02 */ - /* 0x000fd00000000f00 */ - /*0020*/ S2R R4, SR_CTAID.X; /* 0x0000000000047919 */ - /* 0x000e220000002500 */ - /*0030*/ S2R R2, SR_TID.X; /* 0x0000000000027919 */ - /* 0x000e240000002100 */ - /*0040*/ IMAD R4, R4, c[0x0][0x0], R2; /* 0x0000000004047a24 */ - /* 0x001fca00078e0202 */ - /*0050*/ ISETP.GE.AND P0, PT, R4, c[0x0][0x170], PT, !PT; /* 0x00005c0004007a0c */ - /* 0x000fd80003f062f0 */ - /*0060*/ @P0 EXIT; /* 0x000000000000094d */ - /* 0x000fea0003800000 */ - /*0070*/ MOV R2, 0x4; /* 0x0000000400027802 */ - /* 0x000fca0000000f00 */ - /*0080*/ IMAD.WIDE R2, R4, R2, c[0x0][0x168]; /* 0x00005a0004027625 */ - /* 0x000fd400078e0202 */ - /*0090*/ LDG.E.SYS R2, [R2]; /* 0x0000000002027381 */ - /* 0x000e2200001ee900 */ - /*00a0*/ MOV R5, 0x2; /* 0x0000000200057802 */ - /* 0x000fca0000000f00 */ - /*00b0*/ IMAD.WIDE R4, R4, R5, c[0x0][0x160]; /* 0x0000580004047625 */ - /* 0x000fe200078e0205 */ - /*00c0*/ F2F.F16.F32 R6, R2; /* 0x0000000200067304 */ - /* 0x001e320000200800 */ - /*00d0*/ STG.E.U16.SYS [R4], R6; /* 0x0000000604007386 */ - /* 0x0011e2000010e500 */ - /*00e0*/ EXIT; /* 0x000000000000794d */ - /* 0x000fea0003800000 */ - /*00f0*/ BRA 0xf0; /* 0xfffffff000007947 */ - /* 0x000fc0000383ffff */ - ........................................... - - - Function : _Z12wmma_exampleP6__halfS0_Pfiiiff - .headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)" - /*0000*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ; /* 0x000000fffffff389 */ - /* 0x000fe200000e00ff */ - /*0010*/ IMAD.U32 R1, RZ, RZ, c[0x0][0x28]; /* 0x00000a00ff017624 */ - /* 0x000fd000078e00ff */ - /*0020*/ IADD3 R1, R1, -0x8, RZ; /* 0xfffffff801017810 */ - /* 0x000fc80007ffe0ff */ - /*0030*/ IADD3 R2, P0, R1, c[0x0][0x20], RZ; /* 0x0000080001027a10 */ - /* 0x000fca0007f1e0ff */ - /*0040*/ IMAD.X R0, RZ, RZ, c[0x0][0x24], P0; /* 0x00000900ff007624 */ - /* 0x000fd000000e06ff */ - /*0050*/ CS2R.32 R3, SR_CLOCKLO; /* 0x0000000000037805 */ - /* 0x000fd00000005000 */ - /*0060*/ I2F.U32.RP R7, 0x20; /* 0x0000002000077906 */ - /* 0x000e240000209000 */ - /*0070*/ MUFU.RCP R7, R7; /* 0x0000000700077308 */ - /* 0x001e220000001000 */ - /*0080*/ S2R R6, SR_CTAID.X; /* 0x0000000000067919 */ - /* 0x000e620000002500 */ - /*0090*/ S2R R9, SR_TID.X; /* 0x0000000000097919 */ - /* 0x000e620000002100 */ - /*00a0*/ IADD3 R8, R7, 0xffffffe, RZ; /* 0x0ffffffe07087810 */ - /* 0x001fcc0007ffe0ff */ - /*00b0*/ F2I.FTZ.U32.TRUNC.NTZ R5, R8; /* 0x0000000800057305 */ - /* 0x0000a2000021f000 */ - /*00c0*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fe400078e00ff */ - /*00d0*/ IMAD R6, R6, c[0x0][0x0], R9; /* 0x0000000006067a24 */ - /* 0x002fe400078e0209 */ - /*00e0*/ IMAD.U32 R10, R5, -0x20, RZ; /* 0xffffffe0050a7824 */ - /* 0x004fc800078e00ff */ - /*00f0*/ IMAD.WIDE.U32 R4, R5, R10, R4; /* 0x0000000a05047225 */ - /* 0x000fd000078e0004 */ - /*0100*/ IMAD.WIDE.U32 R4, R5, R6, RZ; /* 0x0000000605047225 */ - /* 0x000fcc00078e00ff */ - /*0110*/ IMAD R4, RZ, RZ, -R5; /* 0x000000ffff047224 */ - /* 0x000fc800078e0a05 */ - /*0120*/ IMAD R6, R4, 0x20, R6; /* 0x0000002004067824 */ - /* 0x000fca00078e0206 */ - /*0130*/ ISETP.GE.U32.AND P0, PT, R6.reuse, 0x20, PT, !PT; /* 0x000000200600780c */ - /* 0x040fe20003f060f0 */ - /*0140*/ S2R R28, SR_CTAID.Y; /* 0x00000000001c7919 */ - /* 0x000e220000002600 */ - /*0150*/ S2R R7, SR_TID.Y; /* 0x0000000000077919 */ - /* 0x000e340000002200 */ - /*0160*/ @P0 IADD3 R6, R6, -0x20, RZ; /* 0xffffffe006060810 */ - /* 0x000fc80007ffe0ff */ - /*0170*/ ISETP.GE.U32.AND P1, PT, R6, 0x20, PT, !PT; /* 0x000000200600780c */ - /* 0x000fe40003f260f0 */ - /*0180*/ @P0 IADD3 R5, R5, 0x1, RZ; /* 0x0000000105050810 */ - /* 0x000fe40007ffe0ff */ - /*0190*/ ISETP.LT.AND P0, PT, RZ, c[0x0][0x180], PT, !PT; /* 0x00006000ff007a0c */ - /* 0x000fd00003f012f0 */ - /*01a0*/ @P1 IADD3 R5, R5, 0x1, RZ; /* 0x0000000105051810 */ - /* 0x000fe20007ffe0ff */ - /*01b0*/ IMAD R28, R28, c[0x0][0x4], R7; /* 0x000001001c1c7a24 */ - /* 0x001fc600078e0207 */ - /*01c0*/ SHF.L.U32 R29, R5, 0x4, RZ; /* 0x00000004051d7819 */ - /* 0x000fe200000006ff */ - /*01d0*/ IMAD.U32 R28, R28, 0x10, RZ; /* 0x000000101c1c7824 */ - /* 0x000fc600078e00ff */ - /*01e0*/ ISETP.LT.AND P0, PT, R29, c[0x0][0x178], P0, !PT; /* 0x00005e001d007a0c */ - /* 0x000fc800007012f0 */ - /*01f0*/ ISETP.LT.AND P0, PT, R28, c[0x0][0x17c], P0, !PT; /* 0x00005f001c007a0c */ - /* 0x000fe200007012f0 */ - /*0200*/ BSSY B0, 0x5b0; /* 0x000003a000007945 */ - /* 0x000fe20003800000 */ - /*0210*/ IMAD.U32 R7, RZ, RZ, RZ; /* 0x000000ffff077224 */ - /* 0x000fe200078e00ff */ - /*0220*/ MOV R11, RZ; /* 0x000000ff000b7202 */ - /* 0x000fe20000000f00 */ - /*0230*/ IMAD.U32 R6, RZ, RZ, RZ; /* 0x000000ffff067224 */ - /* 0x000fe400078e00ff */ - /*0240*/ IMAD.U32 R5, RZ, RZ, RZ; /* 0x000000ffff057224 */ - /* 0x000fe400078e00ff */ - /*0250*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fe400078e00ff */ - /*0260*/ IMAD.U32 R10, RZ, RZ, RZ; /* 0x000000ffff0a7224 */ - /* 0x000fc400078e00ff */ - /*0270*/ IMAD.U32 R9, RZ, RZ, RZ; /* 0x000000ffff097224 */ - /* 0x000fe400078e00ff */ - /*0280*/ IMAD.U32 R8, RZ, RZ, RZ; /* 0x000000ffff087224 */ - /* 0x000fe200078e00ff */ - /*0290*/ @!P0 BRA 0x5a0; /* 0x0000030000008947 */ - /* 0x000fee0003800000 */ - /*02a0*/ S2R R6, SR_LANEID; /* 0x0000000000067919 */ - /* 0x000e220000000000 */ - /*02b0*/ IMAD.U32 R10, RZ, RZ, 0x2; /* 0x00000002ff0a7424 */ - /* 0x000fc800078e00ff */ - /*02c0*/ IMAD.WIDE R16, R29, R10, c[0x0][0x160]; /* 0x000058001d107625 */ - /* 0x000fe200078e020a */ - /*02d0*/ SHF.R.U32.HI R4, RZ, 0x2, R6; /* 0x00000002ff047819 */ - /* 0x001fc80000011606 */ - /*02e0*/ LOP3.LUT R5, R4, 0x3, RZ, 0xc0, !PT; /* 0x0000000304057812 */ - /* 0x000fe400078ec0ff */ - /*02f0*/ LOP3.LUT R4, R6, 0x3, RZ, 0xc0, !PT; /* 0x0000000306047812 */ - /* 0x000fe400078ec0ff */ - /*0300*/ LOP3.LUT R7, R5, 0x1, RZ, 0xc0, !PT; /* 0x0000000105077812 */ - /* 0x000fe400078ec0ff */ - /*0310*/ SHF.R.U32.HI R6, RZ, 0x4, R6; /* 0x00000004ff067819 */ - /* 0x000fe40000011606 */ - /*0320*/ SHF.R.U32.HI R5, RZ, 0x1, R5; /* 0x00000001ff057819 */ - /* 0x000fe20000011605 */ - /*0330*/ IMAD R7, R7, 0x8, R4; /* 0x0000000807077824 */ - /* 0x000fe200078e0204 */ - /*0340*/ LOP3.LUT R6, R6, 0x1, RZ, 0xc0, !PT; /* 0x0000000106067812 */ - /* 0x000fc400078ec0ff */ - /*0350*/ LEA R4, R5, R4, 0x3; /* 0x0000000405047211 */ - /* 0x000fe200078e18ff */ - /*0360*/ IMAD.WIDE R12, R28, R10, c[0x0][0x168]; /* 0x00005a001c0c7625 */ - /* 0x000fe400078e020a */ - /*0370*/ IMAD R7, R6.reuse, 0x4, R7; /* 0x0000000406077824 */ - /* 0x040fe400078e0207 */ - /*0380*/ IMAD R4, R6, 0x4, R4; /* 0x0000000406047824 */ - /* 0x000fe400078e0204 */ - /*0390*/ IMAD.U32 R7, R7, 0x2, RZ; /* 0x0000000207077824 */ - /* 0x000fe400078e00ff */ - /*03a0*/ IMAD.U32 R5, R4, 0x2, RZ; /* 0x0000000204057824 */ - /* 0x000fe400078e00ff */ - /*03b0*/ IMAD.WIDE.U32 R16, R7, c[0x0][0x178], R16; /* 0x00005e0007107a25 */ - /* 0x000fc400078e0010 */ - /*03c0*/ IMAD.WIDE.U32 R12, R5, c[0x0][0x180], R12; /* 0x00006000050c7a25 */ - /* 0x000fd000078e000c */ - /*03d0*/ LD.E.128.SYS R24, [R16]; /* 0x0000000010187980 */ - /* 0x000064000010ed00 */ - /*03e0*/ LD.E.128.SYS R20, [R12]; /* 0x000000000c147980 */ - /* 0x000462000010ed00 */ - /*03f0*/ LD.E.128.SYS R16, [R16+0x10]; /* 0x0000001010107980 */ - /* 0x001e22000010ed00 */ - /*0400*/ LD.E.128.SYS R12, [R12+0x10]; /* 0x000000100c0c7980 */ - /* 0x004e22000010ed00 */ - /*0410*/ IMAD.U32 R8, RZ, RZ, RZ; /* 0x000000ffff087224 */ - /* 0x000fe200078e00ff */ - /*0420*/ MOV R9, RZ; /* 0x000000ff00097202 */ - /* 0x000fe20000000f00 */ - /*0430*/ IMAD.U32 R10, RZ, RZ, RZ; /* 0x000000ffff0a7224 */ - /* 0x000fe400078e00ff */ - /*0440*/ IMAD.U32 R11, RZ, RZ, RZ; /* 0x000000ffff0b7224 */ - /* 0x000fe200078e00ff */ - /*0450*/ MOV R7, RZ; /* 0x000000ff00077202 */ - /* 0x000fe20000000f00 */ - /*0460*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fc400078e00ff */ - /*0470*/ IMAD.U32 R5, RZ, RZ, RZ; /* 0x000000ffff057224 */ - /* 0x000fe400078e00ff */ - /*0480*/ IMAD.U32 R6, RZ, RZ, RZ; /* 0x000000ffff067224 */ - /* 0x000fe200078e00ff */ - /*0490*/ WARPSYNC 0xffffffff; /* 0xffffffff00007948 */ - /* 0x000fe20003800000 */ - /*04a0*/ HMMA.884.F32.F32.STEP0 R8, R24.reuse, R20.reuse.T, R8; /* 0x0000001418087236 */ - /* 0x0c22640000005408 */ - /*04b0*/ HMMA.884.F32.F32.STEP1 R10, R24.reuse, R20.reuse.T, R10; /* 0x00000014180a7236 */ - /* 0x0c04a4000000d40a */ - /*04c0*/ HMMA.884.F32.F32.STEP2 R4, R24.reuse, R20.reuse.T, R4; /* 0x0000001418047236 */ - /* 0x0c06e40000015404 */ - /*04d0*/ HMMA.884.F32.F32.STEP3 R6, R24, R20.T, R6; /* 0x0000001418067236 */ - /* 0x000928000001d406 */ - /*04e0*/ HMMA.884.F32.F32.STEP0 R8, R26.reuse, R22.reuse.T, R8; /* 0x000000161a087236 */ - /* 0x0c20240000005408 */ - /*04f0*/ HMMA.884.F32.F32.STEP1 R10, R26.reuse, R22.reuse.T, R10; /* 0x000000161a0a7236 */ - /* 0x0c4264000000d40a */ - /*0500*/ HMMA.884.F32.F32.STEP2 R4, R26.reuse, R22.reuse.T, R4; /* 0x000000161a047236 */ - /* 0x0c84a40000015404 */ - /*0510*/ HMMA.884.F32.F32.STEP3 R6, R26, R22.T, R6; /* 0x000000161a067236 */ - /* 0x0106e8000001d406 */ - /*0520*/ HMMA.884.F32.F32.STEP0 R8, R16.reuse, R12.reuse.T, R8; /* 0x0000000c10087236 */ - /* 0x0c10240000005408 */ - /*0530*/ HMMA.884.F32.F32.STEP1 R10, R16.reuse, R12.reuse.T, R10; /* 0x0000000c100a7236 */ - /* 0x0c2264000000d40a */ - /*0540*/ HMMA.884.F32.F32.STEP2 R4, R16.reuse, R12.reuse.T, R4; /* 0x0000000c10047236 */ - /* 0x0c44a40000015404 */ - /*0550*/ HMMA.884.F32.F32.STEP3 R6, R16, R12.T, R6; /* 0x0000000c10067236 */ - /* 0x0086e8000001d406 */ - /*0560*/ HMMA.884.F32.F32.STEP0 R8, R18.reuse, R14.reuse.T, R8; /* 0x0000000e12087236 */ - /* 0x0c10240000005408 */ - /*0570*/ HMMA.884.F32.F32.STEP1 R10, R18.reuse, R14.reuse.T, R10; /* 0x0000000e120a7236 */ - /* 0x0c2024000000d40a */ - /*0580*/ HMMA.884.F32.F32.STEP2 R4, R18.reuse, R14.reuse.T, R4; /* 0x0000000e12047236 */ - /* 0x0c40240000015404 */ - /*0590*/ HMMA.884.F32.F32.STEP3 R6, R18, R14.T, R6; /* 0x0000000e12067236 */ - /* 0x008034000001d406 */ - /*05a0*/ BSYNC B0; /* 0x0000000000007941 */ - /* 0x001fea0003800000 */ - /*05b0*/ S2R R12, SR_LANEID; /* 0x00000000000c7919 */ - /* 0x000e220000000000 */ - /*05c0*/ IMAD R28, R28, c[0x0][0x178], RZ; /* 0x00005e001c1c7a24 */ - /* 0x000fe200078e02ff */ - /*05d0*/ SHF.R.U32.HI R14, RZ, 0x4, R12; /* 0x00000004ff0e7819 */ - /* 0x001fc4000001160c */ - /*05e0*/ SHF.R.U32.HI R13, RZ, 0x2, R12; /* 0x00000002ff0d7819 */ - /* 0x000fe4000001160c */ - /*05f0*/ LOP3.LUT R12, R12, 0x3, RZ, 0xc0, !PT; /* 0x000000030c0c7812 */ - /* 0x000fe400078ec0ff */ - /*0600*/ LOP3.LUT R14, R14, 0x1, RZ, 0xc0, !PT; /* 0x000000010e0e7812 */ - /* 0x000fe400078ec0ff */ - /*0610*/ LOP3.LUT R13, R13, 0x3, RZ, 0xc0, !PT; /* 0x000000030d0d7812 */ - /* 0x000fc600078ec0ff */ - /*0620*/ IMAD R12, R14, 0x4, R12; /* 0x000000040e0c7824 */ - /* 0x000fe200078e020c */ - /*0630*/ LOP3.LUT R15, R13, 0x1, RZ, 0xc0, !PT; /* 0x000000010d0f7812 */ - /* 0x000fe400078ec0ff */ - /*0640*/ SHF.R.U32.HI R16, RZ, 0x1, R13; /* 0x00000001ff107819 */ - /* 0x000fe4000001160d */ - /*0650*/ LOP3.LUT R13, R12.reuse, 0x5, RZ, 0xc0, !PT; /* 0x000000050c0d7812 */ - /* 0x040fe400078ec0ff */ - /*0660*/ LOP3.LUT R14, R12, 0x2, RZ, 0xc0, !PT; /* 0x000000020c0e7812 */ - /* 0x000fc600078ec0ff */ - /*0670*/ IMAD R12, R15, 0x8, R13; /* 0x000000080f0c7824 */ - /* 0x000fe200078e020d */ - /*0680*/ SHF.R.S32.HI R15, RZ, 0x1f, R29; /* 0x0000001fff0f7819 */ - /* 0x000fe2000001141d */ - /*0690*/ IMAD R14, R16, 0x8, R14; /* 0x00000008100e7824 */ - /* 0x000fe200078e020e */ - /*06a0*/ IADD3 R17, P0, R29, R28, RZ; /* 0x0000001c1d117210 */ - /* 0x000fe20007f1e0ff */ - /*06b0*/ IMAD.U32 R13, RZ, RZ, RZ; /* 0x000000ffff0d7224 */ - /* 0x000fc600078e00ff */ - /*06c0*/ LEA.HI.X.SX32 R15, R28, R15, 0x1, P0; /* 0x0000000f1c0f7211 */ - /* 0x000fe200000f0eff */ - /*06d0*/ IMAD.WIDE.U32 R12, R14, c[0x0][0x178], R12; /* 0x00005e000e0c7a25 */ - /* 0x000fe200078e000c */ - /*06e0*/ LEA R14, P0, R17, c[0x0][0x170], 0x2; /* 0x00005c00110e7a11 */ - /* 0x000fe200078010ff */ - /*06f0*/ IMAD.U32 R18, RZ, RZ, c[0x0][0x178]; /* 0x00005e00ff127624 */ - /* 0x000fc600078e00ff */ - /*0700*/ LEA.HI.X R17, R17, c[0x0][0x174], R15, 0x2, P0; /* 0x00005d0011117a11 */ - /* 0x000fe400000f140f */ - /*0710*/ LEA R16, P0, R12, R14, 0x2; /* 0x0000000e0c107211 */ - /* 0x000fe400078010ff */ - /*0720*/ SHF.L.U32 R14, R18, 0x2, RZ; /* 0x00000002120e7819 */ - /* 0x000fe400000006ff */ - /*0730*/ SHF.R.U32.HI R15, RZ, 0x1e, R18; /* 0x0000001eff0f7819 */ - /* 0x000fe40000011612 */ - /*0740*/ LEA.HI.X R13, R12, R17, R13, 0x2, P0; /* 0x000000110c0d7211 */ - /* 0x000fe400000f140d */ - /*0750*/ LEA R18, P1, R14, R16, 0x2; /* 0x000000100e127211 */ - /* 0x000fc400078210ff */ - /*0760*/ IADD3 R17, P0, R14.reuse, R16, RZ; /* 0x000000100e117210 */ - /* 0x040fe40007f1e0ff */ - /*0770*/ LEA.HI.X R19, R14.reuse, R13, R15, 0x2, P1; /* 0x0000000d0e137211 */ - /* 0x040fe400008f140f */ - /*0780*/ IADD3 R21, P1, R14, R18, RZ; /* 0x000000120e157210 */ - /* 0x000fe20007f3e0ff */ - /*0790*/ IMAD.X R20, R15.reuse, 0x1, R13, P0; /* 0x000000010f147824 */ - /* 0x040fe400000e060d */ - /*07a0*/ IMAD.U32 R12, RZ, RZ, R16; /* 0x000000ffff0c7224 */ - /* 0x000fe400078e0010 */ - /*07b0*/ IMAD.X R22, R15, 0x1, R19, P1; /* 0x000000010f167824 */ - /* 0x000fe400008e0613 */ - /*07c0*/ IMAD.U32 R14, RZ, RZ, R17; /* 0x000000ffff0e7224 */ - /* 0x000fc400078e0011 */ - /*07d0*/ IMAD.U32 R15, RZ, RZ, R20; /* 0x000000ffff0f7224 */ - /* 0x000fe200078e0014 */ - /*07e0*/ MOV R16, R21; /* 0x0000001500107202 */ - /* 0x000fe20000000f00 */ - /*07f0*/ IMAD.U32 R17, RZ, RZ, R22; /* 0x000000ffff117224 */ - /* 0x000fe200078e0016 */ - /*0800*/ ST.E.SYS [R12], R8; /* 0x000000000c007385 */ - /* 0x0001e2000010e908 */ - /*0810*/ ST.E.SYS [R12+0x8], R10; /* 0x000000080c007385 */ - /* 0x0003e8000010e90a */ - /*0820*/ ST.E.SYS [R14], R9; /* 0x000000000e007385 */ - /* 0x0003e2000010e909 */ - /*0830*/ ST.E.SYS [R14+0x8], R11; /* 0x000000080e007385 */ - /* 0x0003e2000010e90b */ - /*0840*/ ST.E.SYS [R18], R4; /* 0x0000000012007385 */ - /* 0x0003e2000010e904 */ - /*0850*/ ST.E.SYS [R18+0x8], R6; /* 0x0000000812007385 */ - /* 0x0003e2000010e906 */ - /*0860*/ ST.E.SYS [R16], R5; /* 0x0000000010007385 */ - /* 0x0003e2000010e905 */ - /*0870*/ ST.E.SYS [R16+0x8], R7; /* 0x0000000810007385 */ - /* 0x0003e2000010e907 */ - /*0880*/ WARPSYNC 0xffffffff; /* 0xffffffff00007948 */ - /* 0x000fe20003800000 */ - /*0890*/ IADD3 R8, R2, -c[0x0][0x20], RZ; /* 0x8000080002087a10 */ - /* 0x001fd00007ffe0ff */ - /*08a0*/ CS2R.32 R4, SR_CLOCKLO; /* 0x0000000000047805 */ - /* 0x002fd00000005000 */ - /*08b0*/ IMAD R3, R4, 0x1, -R3; /* 0x0000000104037824 */ - /* 0x000fd000078e0a03 */ - /*08c0*/ STL [R8], R3; /* 0x0000000308007387 */ - /* 0x0001e20000100800 */ - /*08d0*/ IMAD.U32 R6, RZ, RZ, R2; /* 0x000000ffff067224 */ - /* 0x000fe200078e0002 */ - /*08e0*/ MOV R4, 0x0; /* 0x0000000000047802 */ - /* 0x000fe20000000f00 */ - /*08f0*/ IMAD.U32 R7, RZ, RZ, R0; /* 0x000000ffff077224 */ - /* 0x000fe200078e0000 */ - /*0900*/ MOV R5, 0x0; /* 0x0000000000057802 */ - /* 0x000fe40000000f00 */ - /*0910*/ MOV R20, 0x0; /* 0x0000000000147802 */ - /* 0x000fe40000000f00 */ - /*0920*/ MOV R21, 0x0; /* 0x0000000000157802 */ - /* 0x000fd00000000f00 */ - /*0930*/ CALL.ABS.NOINC 0x0; /* 0x0000000000007943 */ - /* 0x001fea0003c00000 */ - /*0940*/ EXIT; /* 0x000000000000794d */ - /* 0x000fea0003800000 */ - /*0950*/ BRA 0x950; /* 0xfffffff000007947 */ - /* 0x000fc0000383ffff */ - /*0960*/ NOP; /* 0x0000000000007918 */ - /* 0x000fc00000000000 */ - /*0970*/ NOP; /* 0x0000000000007918 */ - /* 0x000fc00000000000 */ - ............................................. - - - -Fatbin ptx code: -================ -arch = sm_70 -code version = [6,0] -producer = cuda -host = linux -compile_size = 64bit -compressed - - - - - - - - -.version 6.0 -.target sm_70 -.address_size 64 - - -.extern .func (.param .b32 func_retval0) vprintf -( -.param .b64 vprintf_param_0, -.param .b64 vprintf_param_1 -) -; -.global .align 16 .b8 $str[9] = {99, 108, 111, 99, 107, 61, 37, 100, 0}; - -.visible .entry _Z12wmma_exampleP6__halfS0_Pfiiiff( -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_0, -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_1, -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_2, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_3, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_4, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_5, -.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_6, -.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_7 -) -{ -.local .align 8 .b8 __local_depot0[8]; -.reg .b64 %SP; -.reg .b64 %SPL; -.reg .pred %p<6>; -.reg .f32 %f<34>; -.reg .b32 %r<38>; -.reg .b64 %rd<18>; - - -mov.u64 %rd17, __local_depot0; -cvta.local.u64 %SP, %rd17; -ld.param.u64 %rd1, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0]; -ld.param.u64 %rd2, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1]; -ld.param.u64 %rd3, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2]; -ld.param.u32 %r4, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3]; -ld.param.u32 %r7, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4]; -ld.param.u32 %r5, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5]; - - mov.u32 %r6, %clock; - - mov.u32 %r8, %ntid.x; -mov.u32 %r9, %ctaid.x; -mov.u32 %r10, %tid.x; -mad.lo.s32 %r11, %r8, %r9, %r10; -mov.u32 %r12, WARP_SZ; -div.u32 %r13, %r11, %r12; -mov.u32 %r14, %ntid.y; -mov.u32 %r15, %ctaid.y; -mov.u32 %r16, %tid.y; -mad.lo.s32 %r17, %r14, %r15, %r16; -shl.b32 %r2, %r13, 4; -shl.b32 %r3, %r17, 4; -setp.lt.s32 %p1, %r2, %r4; -setp.gt.s32 %p2, %r5, 0; -and.pred %p3, %p1, %p2; -setp.lt.s32 %p4, %r3, %r7; -and.pred %p5, %p3, %p4; -mov.f32 %f26, 0f00000000; -mov.f32 %f27, %f26; -mov.f32 %f28, %f26; -mov.f32 %f29, %f26; -mov.f32 %f30, %f26; -mov.f32 %f31, %f26; -mov.f32 %f32, %f26; -mov.f32 %f33, %f26; -@!%p5 bra BB0_2; -bra.uni BB0_1; - -BB0_1: -mul.wide.s32 %rd4, %r2, 2; -add.s64 %rd5, %rd1, %rd4; -wmma.load.a.sync.row.m16n16k16.f16 {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, [%rd5], %r4; -mul.wide.s32 %rd6, %r3, 2; -add.s64 %rd7, %rd2, %rd6; -wmma.load.b.sync.col.m16n16k16.f16 {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, [%rd7], %r5; -mov.f32 %f25, 0f00000000; -wmma.mma.sync.row.col.m16n16k16.f32.f32 {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, {%f25, %f25, %f25, %f25, %f25, %f25, %f25, %f25}; - -BB0_2: -add.u64 %rd8, %SP, 0; -cvta.to.local.u64 %rd9, %rd8; -mul.lo.s32 %r35, %r3, %r4; -cvt.s64.s32 %rd10, %r35; -cvt.s64.s32 %rd11, %r2; -add.s64 %rd12, %rd10, %rd11; -shl.b64 %rd13, %rd12, 2; -add.s64 %rd14, %rd3, %rd13; -wmma.store.d.sync.col.m16n16k16.f32 [%rd14], {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, %r4; - - mov.u32 %r34, %clock; - - sub.s32 %r36, %r34, %r6; -st.local.u32 [%rd9], %r36; -mov.u64 %rd15, $str; -cvta.global.u64 %rd16, %rd15; - - { -.reg .b32 temp_param_reg; - - .param .b64 param0; -st.param.b64 [param0+0], %rd16; -.param .b64 param1; -st.param.b64 [param1+0], %rd8; -.param .b32 retval0; -call.uni (retval0), -vprintf, -( -param0, -param1 -); -ld.param.b32 %r37, [retval0+0]; - - - } - ret; -} - - -.visible .entry _Z17convertFp32ToFp16P6__halfPfi( -.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_0, -.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_1, -.param .u32 _Z17convertFp32ToFp16P6__halfPfi_param_2 -) -{ -.reg .pred %p<2>; -.reg .b16 %rs<2>; -.reg .f32 %f<2>; -.reg .b32 %r<6>; -.reg .b64 %rd<9>; - - -ld.param.u64 %rd1, [_Z17convertFp32ToFp16P6__halfPfi_param_0]; -ld.param.u64 %rd2, [_Z17convertFp32ToFp16P6__halfPfi_param_1]; -ld.param.u32 %r2, [_Z17convertFp32ToFp16P6__halfPfi_param_2]; -mov.u32 %r3, %ntid.x; -mov.u32 %r4, %ctaid.x; -mov.u32 %r5, %tid.x; -mad.lo.s32 %r1, %r4, %r3, %r5; -setp.ge.s32 %p1, %r1, %r2; -@%p1 bra BB1_2; - -cvta.to.global.u64 %rd3, %rd2; -mul.wide.s32 %rd4, %r1, 4; -add.s64 %rd5, %rd3, %rd4; -ld.global.f32 %f1, [%rd5]; - - { cvt.rn.f16.f32 %rs1, %f1;} - - - cvta.to.global.u64 %rd6, %rd1; -mul.wide.s32 %rd7, %r1, 2; -add.s64 %rd8, %rd6, %rd7; -st.global.u16 [%rd8], %rs1; - -BB1_2: -ret; -} - - diff --git a/cuda-kernels/_cuobjdump_complete_output_rndQyq b/cuda-kernels/_cuobjdump_complete_output_rndQyq deleted file mode 100644 index 36999c0..0000000 --- a/cuda-kernels/_cuobjdump_complete_output_rndQyq +++ /dev/null @@ -1,1055 +0,0 @@ - -Fatbin elf code: -================ -arch = sm_70 -code version = [1,7] -producer = -host = linux -compile_size = 64bit - -64bit elf: type=2, abi=7, sm=70, toolkit=90, flags = 0x460546 -Sections: -Index Offset Size ES Align Type Flags Link Info Name - 1 40 32 0 1 STRTAB 0 0 0 .shstrtab - 2 72 32 0 1 STRTAB 0 0 0 .strtab - 3 a8 18 18 8 SYMTAB 0 2 0 .symtab - -.section .strtab - -.section .shstrtab - -.section .symtab - index value size info other shndx name - 0 0 0 0 0 0 (null) - - code for sm_70 - -Fatbin elf code: -================ -arch = sm_70 -code version = [1,7] -producer = cuda -host = linux -compile_size = 64bit - -64bit elf: type=2, abi=7, sm=70, toolkit=90, flags = 0x460546 -Sections: -Index Offset Size ES Align Type Flags Link Info Name - 1 40 21b 0 1 STRTAB 0 0 0 .shstrtab - 2 25b 273 0 1 STRTAB 0 0 0 .strtab - 3 4d0 108 18 8 SYMTAB 0 2 7 .symtab - 4 5d8 e0 0 1 PROGBITS 0 0 0 .debug_frame - 5 6b8 48 0 4 CUDA_INFO 0 3 0 .nv.info - 6 700 50 0 4 CUDA_INFO 0 3 d .nv.info._Z17convertFp32ToFp16P6__halfPfi - 7 750 ac 0 4 CUDA_INFO 0 3 e .nv.info._Z12wmma_exampleP6__halfS0_Pfiiiff - 8 800 30 10 8 REL 0 3 e .rel.text._Z12wmma_exampleP6__halfS0_Pfiiiff - 9 830 30 18 8 RELA 0 3 e .rela.text._Z12wmma_exampleP6__halfS0_Pfiiiff - a 860 20 10 8 REL 0 3 4 .rel.debug_frame - b 880 174 0 4 PROGBITS 2 0 d .nv.constant0._Z17convertFp32ToFp16P6__halfPfi - c 9f4 18c 0 4 PROGBITS 2 0 e .nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff - d b80 100 0 80 PROGBITS 6 3 9000008 .text._Z17convertFp32ToFp16P6__halfPfi - e c80 980 0 80 PROGBITS 6 3 20000009 .text._Z12wmma_exampleP6__halfS0_Pfiiiff - f 1600 9 0 10 PROGBITS 3 0 0 .nv.global.init - -.section .strtab - -.section .shstrtab - -.section .symtab - index value size info other shndx name - 0 0 0 0 0 0 (null) - 1 0 0 3 0 d .text._Z17convertFp32ToFp16P6__halfPfi - 2 0 0 3 0 f .nv.global.init - 3 0 9 1 0 f $str - 4 0 0 3 0 b .nv.constant0._Z17convertFp32ToFp16P6__halfPfi - 5 0 0 3 0 e .text._Z12wmma_exampleP6__halfS0_Pfiiiff - 6 0 0 3 0 c .nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff - 7 0 0 3 0 4 .debug_frame - 8 0 256 12 10 d _Z17convertFp32ToFp16P6__halfPfi - 9 0 2432 12 10 e _Z12wmma_exampleP6__halfS0_Pfiiiff - 10 0 0 12 0 0 vprintf - - -.nv.constant0._Z17convertFp32ToFp16P6__halfPfi -0x00000000 0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 - - - -.nv.constant0._Z12wmma_exampleP6__halfS0_Pfiiiff -0x00000000 0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 0x00000000 0x00000000 -0x00000000 0x00000000 - - -.nv.global.init -0x636f6c63 0x64253d6b 0 - - -.nv.info - <0x1> - Attribute: EIATTR_MAX_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x9 0x0 - <0x2> - Attribute: EIATTR_MIN_STACK_SIZE - Format: EIFMT_SVAL - Value: function: _Z12wmma_exampleP6__halfS0_Pfiiiff(0x9) min stack size: 0x8 - <0x3> - Attribute: EIATTR_FRAME_SIZE - Format: EIFMT_SVAL - Value: function: _Z12wmma_exampleP6__halfS0_Pfiiiff(0x9) frame size: 0x8 - <0x4> - Attribute: EIATTR_MAX_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x8 0x0 - <0x5> - Attribute: EIATTR_MIN_STACK_SIZE - Format: EIFMT_SVAL - Value: function: _Z17convertFp32ToFp16P6__halfPfi(0x8) min stack size: 0x0 - <0x6> - Attribute: EIATTR_FRAME_SIZE - Format: EIFMT_SVAL - Value: function: _Z17convertFp32ToFp16P6__halfPfi(0x8) frame size: 0x0 - - -.nv.info._Z17convertFp32ToFp16P6__halfPfi - <0x1> - Attribute: EIATTR_PARAM_CBANK - Format: EIFMT_SVAL - Value: 0x4 0x140160 - <0x2> - Attribute: EIATTR_CBANK_PARAM_SIZE - Format: EIFMT_HVAL - Value: 0x14 - <0x3> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x2 Offset : 0x10 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x4> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x1 Offset : 0x8 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x5> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x0 Offset : 0x0 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x6> - Attribute: EIATTR_MAXREG_COUNT - Format: EIFMT_HVAL - Value: 0xff - <0x7> - Attribute: EIATTR_EXIT_INSTR_OFFSETS - Format: EIFMT_SVAL - Value: 0x60 0xe0 - - -.nv.info._Z12wmma_exampleP6__halfS0_Pfiiiff - <0x1> - Attribute: EIATTR_PARAM_CBANK - Format: EIFMT_SVAL - Value: 0x6 0x2c0160 - <0x2> - Attribute: EIATTR_CBANK_PARAM_SIZE - Format: EIFMT_HVAL - Value: 0x2c - <0x3> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x7 Offset : 0x28 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x4> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x6 Offset : 0x24 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x5> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x5 Offset : 0x20 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x6> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x4 Offset : 0x1c Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x7> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x3 Offset : 0x18 Size : 0x4 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x8> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x2 Offset : 0x10 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x9> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x1 Offset : 0x8 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x10> - Attribute: EIATTR_KPARAM_INFO - Format: EIFMT_SVAL - Value: Index : 0x0 Ordinal : 0x0 Offset : 0x0 Size : 0x8 - Pointee's logAlignment : 0x0 Space : 0x0 cbank : 0x1f Parameter Space : CBANK - <0x11> - Attribute: EIATTR_MAXREG_COUNT - Format: EIFMT_HVAL - Value: 0xff - <0x12> - Attribute: EIATTR_EXIT_INSTR_OFFSETS - Format: EIFMT_SVAL - Value: 0x940 - <0x13> - Attribute: EIATTR_EXTERNS - Format: EIFMT_SVAL - Value: externs: vprintf(0xa) - <0x14> - Attribute: EIATTR_CRS_STACK_SIZE - Format: EIFMT_SVAL - Value: 0x0 - - -.text._Z17convertFp32ToFp16P6__halfPfi -bar = 0 reg = 9 lmem=0 smem=0 -0xfffff389 0x000000ff 0x000e00ff 0x000fe200 -0x00017a02 0x00000a00 0x00000f00 0x000fd000 -0x00047919 0x00000000 0x00002500 0x000e2200 -0x00027919 0x00000000 0x00002100 0x000e2400 -0x04047a24 0x00000000 0x078e0202 0x001fca00 -0x04007a0c 0x00005c00 0x03f062f0 0x000fd800 -0x0000094d 0x00000000 0x03800000 0x000fea00 -0x00027802 0x00000004 0x00000f00 0x000fca00 -0x04027625 0x00005a00 0x078e0202 0x000fd400 -0x02027381 0x00000000 0x001ee900 0x000e2200 -0x00057802 0x00000002 0x00000f00 0x000fca00 -0x04047625 0x00005800 0x078e0205 0x000fe200 -0x00067304 0x00000002 0x00200800 0x001e3200 -0x04007386 0x00000006 0x0010e500 0x0011e200 -0x0000794d 0x00000000 0x03800000 0x000fea00 -0x00007947 0xfffffff0 0x0383ffff 0x000fc000 - - - -.text._Z12wmma_exampleP6__halfS0_Pfiiiff -bar = 0 reg = 32 lmem=0 smem=0 -0xfffff389 0x000000ff 0x000e00ff 0x000fe200 -0xff017624 0x00000a00 0x078e00ff 0x000fd000 -0x01017810 0xfffffff8 0x07ffe0ff 0x000fc800 -0x01027a10 0x00000800 0x07f1e0ff 0x000fca00 -0xff007624 0x00000900 0x000e06ff 0x000fd000 -0x00037805 0x00000000 0x00005000 0x000fd000 -0x00077906 0x00000020 0x00209000 0x000e2400 -0x00077308 0x00000007 0x00001000 0x001e2200 -0x00067919 0x00000000 0x00002500 0x000e6200 -0x00097919 0x00000000 0x00002100 0x000e6200 -0x07087810 0x0ffffffe 0x07ffe0ff 0x001fcc00 -0x00057305 0x00000008 0x0021f000 0x0000a200 -0xff047224 0x000000ff 0x078e00ff 0x000fe400 -0x06067a24 0x00000000 0x078e0209 0x002fe400 -0x050a7824 0xffffffe0 0x078e00ff 0x004fc800 -0x05047225 0x0000000a 0x078e0004 0x000fd000 -0x05047225 0x00000006 0x078e00ff 0x000fcc00 -0xff047224 0x000000ff 0x078e0a05 0x000fc800 -0x04067824 0x00000020 0x078e0206 0x000fca00 -0x0600780c 0x00000020 0x03f060f0 0x040fe200 -0x001c7919 0x00000000 0x00002600 0x000e2200 -0x00077919 0x00000000 0x00002200 0x000e3400 -0x06060810 0xffffffe0 0x07ffe0ff 0x000fc800 -0x0600780c 0x00000020 0x03f260f0 0x000fe400 -0x05050810 0x00000001 0x07ffe0ff 0x000fe400 -0xff007a0c 0x00006000 0x03f012f0 0x000fd000 -0x05051810 0x00000001 0x07ffe0ff 0x000fe200 -0x1c1c7a24 0x00000100 0x078e0207 0x001fc600 -0x051d7819 0x00000004 0x000006ff 0x000fe200 -0x1c1c7824 0x00000010 0x078e00ff 0x000fc600 -0x1d007a0c 0x00005e00 0x007012f0 0x000fc800 -0x1c007a0c 0x00005f00 0x007012f0 0x000fe200 -0x00007945 0x000003a0 0x03800000 0x000fe200 -0xff077224 0x000000ff 0x078e00ff 0x000fe200 -0x000b7202 0x000000ff 0x00000f00 0x000fe200 -0xff067224 0x000000ff 0x078e00ff 0x000fe400 -0xff057224 0x000000ff 0x078e00ff 0x000fe400 -0xff047224 0x000000ff 0x078e00ff 0x000fe400 -0xff0a7224 0x000000ff 0x078e00ff 0x000fc400 -0xff097224 0x000000ff 0x078e00ff 0x000fe400 -0xff087224 0x000000ff 0x078e00ff 0x000fe200 -0x00008947 0x00000300 0x03800000 0x000fee00 -0x00067919 0x00000000 0x00000000 0x000e2200 -0xff0a7424 0x00000002 0x078e00ff 0x000fc800 -0x1d107625 0x00005800 0x078e020a 0x000fe200 -0xff047819 0x00000002 0x00011606 0x001fc800 -0x04057812 0x00000003 0x078ec0ff 0x000fe400 -0x06047812 0x00000003 0x078ec0ff 0x000fe400 -0x05077812 0x00000001 0x078ec0ff 0x000fe400 -0xff067819 0x00000004 0x00011606 0x000fe400 -0xff057819 0x00000001 0x00011605 0x000fe200 -0x07077824 0x00000008 0x078e0204 0x000fe200 -0x06067812 0x00000001 0x078ec0ff 0x000fc400 -0x05047211 0x00000004 0x078e18ff 0x000fe200 -0x1c0c7625 0x00005a00 0x078e020a 0x000fe400 -0x06077824 0x00000004 0x078e0207 0x040fe400 -0x06047824 0x00000004 0x078e0204 0x000fe400 -0x07077824 0x00000002 0x078e00ff 0x000fe400 -0x04057824 0x00000002 0x078e00ff 0x000fe400 -0x07107a25 0x00005e00 0x078e0010 0x000fc400 -0x050c7a25 0x00006000 0x078e000c 0x000fd000 -0x10187980 0x00000000 0x0010ed00 0x00006400 -0x0c147980 0x00000000 0x0010ed00 0x00046200 -0x10107980 0x00000010 0x0010ed00 0x001e2200 -0x0c0c7980 0x00000010 0x0010ed00 0x004e2200 -0xff087224 0x000000ff 0x078e00ff 0x000fe200 -0x00097202 0x000000ff 0x00000f00 0x000fe200 -0xff0a7224 0x000000ff 0x078e00ff 0x000fe400 -0xff0b7224 0x000000ff 0x078e00ff 0x000fe200 -0x00077202 0x000000ff 0x00000f00 0x000fe200 -0xff047224 0x000000ff 0x078e00ff 0x000fc400 -0xff057224 0x000000ff 0x078e00ff 0x000fe400 -0xff067224 0x000000ff 0x078e00ff 0x000fe200 -0x00007948 0xffffffff 0x03800000 0x000fe200 -0x18087236 0x00000014 0x00005408 0x0c226400 -0x180a7236 0x00000014 0x0000d40a 0x0c04a400 -0x18047236 0x00000014 0x00015404 0x0c06e400 -0x18067236 0x00000014 0x0001d406 0x00092800 -0x1a087236 0x00000016 0x00005408 0x0c202400 -0x1a0a7236 0x00000016 0x0000d40a 0x0c426400 -0x1a047236 0x00000016 0x00015404 0x0c84a400 -0x1a067236 0x00000016 0x0001d406 0x0106e800 -0x10087236 0x0000000c 0x00005408 0x0c102400 -0x100a7236 0x0000000c 0x0000d40a 0x0c226400 -0x10047236 0x0000000c 0x00015404 0x0c44a400 -0x10067236 0x0000000c 0x0001d406 0x0086e800 -0x12087236 0x0000000e 0x00005408 0x0c102400 -0x120a7236 0x0000000e 0x0000d40a 0x0c202400 -0x12047236 0x0000000e 0x00015404 0x0c402400 -0x12067236 0x0000000e 0x0001d406 0x00803400 -0x00007941 0x00000000 0x03800000 0x001fea00 -0x000c7919 0x00000000 0x00000000 0x000e2200 -0x1c1c7a24 0x00005e00 0x078e02ff 0x000fe200 -0xff0e7819 0x00000004 0x0001160c 0x001fc400 -0xff0d7819 0x00000002 0x0001160c 0x000fe400 -0x0c0c7812 0x00000003 0x078ec0ff 0x000fe400 -0x0e0e7812 0x00000001 0x078ec0ff 0x000fe400 -0x0d0d7812 0x00000003 0x078ec0ff 0x000fc600 -0x0e0c7824 0x00000004 0x078e020c 0x000fe200 -0x0d0f7812 0x00000001 0x078ec0ff 0x000fe400 -0xff107819 0x00000001 0x0001160d 0x000fe400 -0x0c0d7812 0x00000005 0x078ec0ff 0x040fe400 -0x0c0e7812 0x00000002 0x078ec0ff 0x000fc600 -0x0f0c7824 0x00000008 0x078e020d 0x000fe200 -0xff0f7819 0x0000001f 0x0001141d 0x000fe200 -0x100e7824 0x00000008 0x078e020e 0x000fe200 -0x1d117210 0x0000001c 0x07f1e0ff 0x000fe200 -0xff0d7224 0x000000ff 0x078e00ff 0x000fc600 -0x1c0f7211 0x0000000f 0x000f0eff 0x000fe200 -0x0e0c7a25 0x00005e00 0x078e000c 0x000fe200 -0x110e7a11 0x00005c00 0x078010ff 0x000fe200 -0xff127624 0x00005e00 0x078e00ff 0x000fc600 -0x11117a11 0x00005d00 0x000f140f 0x000fe400 -0x0c107211 0x0000000e 0x078010ff 0x000fe400 -0x120e7819 0x00000002 0x000006ff 0x000fe400 -0xff0f7819 0x0000001e 0x00011612 0x000fe400 -0x0c0d7211 0x00000011 0x000f140d 0x000fe400 -0x0e127211 0x00000010 0x078210ff 0x000fc400 -0x0e117210 0x00000010 0x07f1e0ff 0x040fe400 -0x0e137211 0x0000000d 0x008f140f 0x040fe400 -0x0e157210 0x00000012 0x07f3e0ff 0x000fe200 -0x0f147824 0x00000001 0x000e060d 0x040fe400 -0xff0c7224 0x000000ff 0x078e0010 0x000fe400 -0x0f167824 0x00000001 0x008e0613 0x000fe400 -0xff0e7224 0x000000ff 0x078e0011 0x000fc400 -0xff0f7224 0x000000ff 0x078e0014 0x000fe200 -0x00107202 0x00000015 0x00000f00 0x000fe200 -0xff117224 0x000000ff 0x078e0016 0x000fe200 -0x0c007385 0x00000000 0x0010e908 0x0001e200 -0x0c007385 0x00000008 0x0010e90a 0x0003e800 -0x0e007385 0x00000000 0x0010e909 0x0003e200 -0x0e007385 0x00000008 0x0010e90b 0x0003e200 -0x12007385 0x00000000 0x0010e904 0x0003e200 -0x12007385 0x00000008 0x0010e906 0x0003e200 -0x10007385 0x00000000 0x0010e905 0x0003e200 -0x10007385 0x00000008 0x0010e907 0x0003e200 -0x00007948 0xffffffff 0x03800000 0x000fe200 -0x02087a10 0x80000800 0x07ffe0ff 0x001fd000 -0x00047805 0x00000000 0x00005000 0x002fd000 -0x04037824 0x00000001 0x078e0a03 0x000fd000 -0x08007387 0x00000003 0x00100800 0x0001e200 -0xff067224 0x000000ff 0x078e0002 0x000fe200 -0x00047802 0x00000000 0x00000f00 0x000fe200 -0xff077224 0x000000ff 0x078e0000 0x000fe200 -0x00057802 0x00000000 0x00000f00 0x000fe400 -0x00147802 0x00000000 0x00000f00 0x000fe400 -0x00157802 0x00000000 0x00000f00 0x000fd000 -0x00007943 0x00000000 0x03c00000 0x001fea00 -0x0000794d 0x00000000 0x03800000 0x000fea00 -0x00007947 0xfffffff0 0x0383ffff 0x000fc000 -0x00007918 0x00000000 0x00000000 0x000fc000 -0x00007918 0x00000000 0x00000000 0x000fc000 - - -.section .rel.text._Z12wmma_exampleP6__halfS0_Pfiiiff REL -2272 $str R_CUDA_ABS32_LO_32 -2304 $str R_CUDA_ABS32_HI_32 -2352 vprintf R_CUDA_ABS47_34 - -.section .rela.text._Z12wmma_exampleP6__halfS0_Pfiiiff RELA -2320 _Z12wmma_exampleP6__halfS0_Pfiiiff R_CUDA_ABS32_LO_32 2368 -2336 _Z12wmma_exampleP6__halfS0_Pfiiiff R_CUDA_ABS32_HI_32 2368 - -.section .debug_frame -decodeDebugFrame, frameBuf 0xffffffff, total_length 224 -CIE length 40, cie_id -1 -version 3 -augmentation slen 1 -augmentation -code_align_factor slen 1 -data_align_factor slen 1 - Debug Frame Common Information Entry - length: 40 - CIE_id : -1 - version: 3 - augmentation: - code align factor: 4 - data align factor: -4 - return address register 16777215 - initial instructions: 23 bytes, ptr = 0x8080810c, frameBuf = 0xffffffff - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_same_value R255 - DW_CFA_same_value R1 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - Debug Frame Description Entry - length: 48 - CIE_pointer: 0 - initial_location: 0x0 - address_range: 0x100 - instructions: 24 bytes - DW_CFA_advance_loc4 delta 4 - DW_CFA_advance_loc4 delta 0 - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_advance_loc4 delta 52 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop -CIE length 40, cie_id -1 -version 3 -augmentation slen 1 -augmentation -code_align_factor slen 1 -data_align_factor slen 1 - Debug Frame Common Information Entry - length: 40 - CIE_id : -1 - version: 3 - augmentation: - code align factor: 4 - data align factor: -4 - return address register 16777215 - initial instructions: 23 bytes, ptr = 0x8080810c, frameBuf = 0xffffffff - DW_CFA_def_cfa register R1, offset 0 - DW_CFA_same_value R255 - DW_CFA_same_value R1 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - Debug Frame Description Entry - length: 48 - CIE_pointer: 0 - initial_location: 0x0 - address_range: 0x970 - instructions: 24 bytes - DW_CFA_advance_loc4 delta 4 - DW_CFA_advance_loc4 delta 2 - DW_CFA_def_cfa register R1, offset 8 - DW_CFA_advance_loc4 delta 586 - DW_CFA_nop - DW_CFA_nop - DW_CFA_nop - -.section .rel.debug_frame REL -72 _Z17convertFp32ToFp16P6__halfPfi R_NV_64 -184 _Z12wmma_exampleP6__halfS0_Pfiiiff R_NV_64 - - code for sm_70 - Function : _Z17convertFp32ToFp16P6__halfPfi - .headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)" - /*0000*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ; /* 0x000000fffffff389 */ - /* 0x000fe200000e00ff */ - /*0010*/ MOV R1, c[0x0][0x28]; /* 0x00000a0000017a02 */ - /* 0x000fd00000000f00 */ - /*0020*/ S2R R4, SR_CTAID.X; /* 0x0000000000047919 */ - /* 0x000e220000002500 */ - /*0030*/ S2R R2, SR_TID.X; /* 0x0000000000027919 */ - /* 0x000e240000002100 */ - /*0040*/ IMAD R4, R4, c[0x0][0x0], R2; /* 0x0000000004047a24 */ - /* 0x001fca00078e0202 */ - /*0050*/ ISETP.GE.AND P0, PT, R4, c[0x0][0x170], PT, !PT; /* 0x00005c0004007a0c */ - /* 0x000fd80003f062f0 */ - /*0060*/ @P0 EXIT; /* 0x000000000000094d */ - /* 0x000fea0003800000 */ - /*0070*/ MOV R2, 0x4; /* 0x0000000400027802 */ - /* 0x000fca0000000f00 */ - /*0080*/ IMAD.WIDE R2, R4, R2, c[0x0][0x168]; /* 0x00005a0004027625 */ - /* 0x000fd400078e0202 */ - /*0090*/ LDG.E.SYS R2, [R2]; /* 0x0000000002027381 */ - /* 0x000e2200001ee900 */ - /*00a0*/ MOV R5, 0x2; /* 0x0000000200057802 */ - /* 0x000fca0000000f00 */ - /*00b0*/ IMAD.WIDE R4, R4, R5, c[0x0][0x160]; /* 0x0000580004047625 */ - /* 0x000fe200078e0205 */ - /*00c0*/ F2F.F16.F32 R6, R2; /* 0x0000000200067304 */ - /* 0x001e320000200800 */ - /*00d0*/ STG.E.U16.SYS [R4], R6; /* 0x0000000604007386 */ - /* 0x0011e2000010e500 */ - /*00e0*/ EXIT; /* 0x000000000000794d */ - /* 0x000fea0003800000 */ - /*00f0*/ BRA 0xf0; /* 0xfffffff000007947 */ - /* 0x000fc0000383ffff */ - ........................................... - - - Function : _Z12wmma_exampleP6__halfS0_Pfiiiff - .headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)" - /*0000*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ; /* 0x000000fffffff389 */ - /* 0x000fe200000e00ff */ - /*0010*/ IMAD.U32 R1, RZ, RZ, c[0x0][0x28]; /* 0x00000a00ff017624 */ - /* 0x000fd000078e00ff */ - /*0020*/ IADD3 R1, R1, -0x8, RZ; /* 0xfffffff801017810 */ - /* 0x000fc80007ffe0ff */ - /*0030*/ IADD3 R2, P0, R1, c[0x0][0x20], RZ; /* 0x0000080001027a10 */ - /* 0x000fca0007f1e0ff */ - /*0040*/ IMAD.X R0, RZ, RZ, c[0x0][0x24], P0; /* 0x00000900ff007624 */ - /* 0x000fd000000e06ff */ - /*0050*/ CS2R.32 R3, SR_CLOCKLO; /* 0x0000000000037805 */ - /* 0x000fd00000005000 */ - /*0060*/ I2F.U32.RP R7, 0x20; /* 0x0000002000077906 */ - /* 0x000e240000209000 */ - /*0070*/ MUFU.RCP R7, R7; /* 0x0000000700077308 */ - /* 0x001e220000001000 */ - /*0080*/ S2R R6, SR_CTAID.X; /* 0x0000000000067919 */ - /* 0x000e620000002500 */ - /*0090*/ S2R R9, SR_TID.X; /* 0x0000000000097919 */ - /* 0x000e620000002100 */ - /*00a0*/ IADD3 R8, R7, 0xffffffe, RZ; /* 0x0ffffffe07087810 */ - /* 0x001fcc0007ffe0ff */ - /*00b0*/ F2I.FTZ.U32.TRUNC.NTZ R5, R8; /* 0x0000000800057305 */ - /* 0x0000a2000021f000 */ - /*00c0*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fe400078e00ff */ - /*00d0*/ IMAD R6, R6, c[0x0][0x0], R9; /* 0x0000000006067a24 */ - /* 0x002fe400078e0209 */ - /*00e0*/ IMAD.U32 R10, R5, -0x20, RZ; /* 0xffffffe0050a7824 */ - /* 0x004fc800078e00ff */ - /*00f0*/ IMAD.WIDE.U32 R4, R5, R10, R4; /* 0x0000000a05047225 */ - /* 0x000fd000078e0004 */ - /*0100*/ IMAD.WIDE.U32 R4, R5, R6, RZ; /* 0x0000000605047225 */ - /* 0x000fcc00078e00ff */ - /*0110*/ IMAD R4, RZ, RZ, -R5; /* 0x000000ffff047224 */ - /* 0x000fc800078e0a05 */ - /*0120*/ IMAD R6, R4, 0x20, R6; /* 0x0000002004067824 */ - /* 0x000fca00078e0206 */ - /*0130*/ ISETP.GE.U32.AND P0, PT, R6.reuse, 0x20, PT, !PT; /* 0x000000200600780c */ - /* 0x040fe20003f060f0 */ - /*0140*/ S2R R28, SR_CTAID.Y; /* 0x00000000001c7919 */ - /* 0x000e220000002600 */ - /*0150*/ S2R R7, SR_TID.Y; /* 0x0000000000077919 */ - /* 0x000e340000002200 */ - /*0160*/ @P0 IADD3 R6, R6, -0x20, RZ; /* 0xffffffe006060810 */ - /* 0x000fc80007ffe0ff */ - /*0170*/ ISETP.GE.U32.AND P1, PT, R6, 0x20, PT, !PT; /* 0x000000200600780c */ - /* 0x000fe40003f260f0 */ - /*0180*/ @P0 IADD3 R5, R5, 0x1, RZ; /* 0x0000000105050810 */ - /* 0x000fe40007ffe0ff */ - /*0190*/ ISETP.LT.AND P0, PT, RZ, c[0x0][0x180], PT, !PT; /* 0x00006000ff007a0c */ - /* 0x000fd00003f012f0 */ - /*01a0*/ @P1 IADD3 R5, R5, 0x1, RZ; /* 0x0000000105051810 */ - /* 0x000fe20007ffe0ff */ - /*01b0*/ IMAD R28, R28, c[0x0][0x4], R7; /* 0x000001001c1c7a24 */ - /* 0x001fc600078e0207 */ - /*01c0*/ SHF.L.U32 R29, R5, 0x4, RZ; /* 0x00000004051d7819 */ - /* 0x000fe200000006ff */ - /*01d0*/ IMAD.U32 R28, R28, 0x10, RZ; /* 0x000000101c1c7824 */ - /* 0x000fc600078e00ff */ - /*01e0*/ ISETP.LT.AND P0, PT, R29, c[0x0][0x178], P0, !PT; /* 0x00005e001d007a0c */ - /* 0x000fc800007012f0 */ - /*01f0*/ ISETP.LT.AND P0, PT, R28, c[0x0][0x17c], P0, !PT; /* 0x00005f001c007a0c */ - /* 0x000fe200007012f0 */ - /*0200*/ BSSY B0, 0x5b0; /* 0x000003a000007945 */ - /* 0x000fe20003800000 */ - /*0210*/ IMAD.U32 R7, RZ, RZ, RZ; /* 0x000000ffff077224 */ - /* 0x000fe200078e00ff */ - /*0220*/ MOV R11, RZ; /* 0x000000ff000b7202 */ - /* 0x000fe20000000f00 */ - /*0230*/ IMAD.U32 R6, RZ, RZ, RZ; /* 0x000000ffff067224 */ - /* 0x000fe400078e00ff */ - /*0240*/ IMAD.U32 R5, RZ, RZ, RZ; /* 0x000000ffff057224 */ - /* 0x000fe400078e00ff */ - /*0250*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fe400078e00ff */ - /*0260*/ IMAD.U32 R10, RZ, RZ, RZ; /* 0x000000ffff0a7224 */ - /* 0x000fc400078e00ff */ - /*0270*/ IMAD.U32 R9, RZ, RZ, RZ; /* 0x000000ffff097224 */ - /* 0x000fe400078e00ff */ - /*0280*/ IMAD.U32 R8, RZ, RZ, RZ; /* 0x000000ffff087224 */ - /* 0x000fe200078e00ff */ - /*0290*/ @!P0 BRA 0x5a0; /* 0x0000030000008947 */ - /* 0x000fee0003800000 */ - /*02a0*/ S2R R6, SR_LANEID; /* 0x0000000000067919 */ - /* 0x000e220000000000 */ - /*02b0*/ IMAD.U32 R10, RZ, RZ, 0x2; /* 0x00000002ff0a7424 */ - /* 0x000fc800078e00ff */ - /*02c0*/ IMAD.WIDE R16, R29, R10, c[0x0][0x160]; /* 0x000058001d107625 */ - /* 0x000fe200078e020a */ - /*02d0*/ SHF.R.U32.HI R4, RZ, 0x2, R6; /* 0x00000002ff047819 */ - /* 0x001fc80000011606 */ - /*02e0*/ LOP3.LUT R5, R4, 0x3, RZ, 0xc0, !PT; /* 0x0000000304057812 */ - /* 0x000fe400078ec0ff */ - /*02f0*/ LOP3.LUT R4, R6, 0x3, RZ, 0xc0, !PT; /* 0x0000000306047812 */ - /* 0x000fe400078ec0ff */ - /*0300*/ LOP3.LUT R7, R5, 0x1, RZ, 0xc0, !PT; /* 0x0000000105077812 */ - /* 0x000fe400078ec0ff */ - /*0310*/ SHF.R.U32.HI R6, RZ, 0x4, R6; /* 0x00000004ff067819 */ - /* 0x000fe40000011606 */ - /*0320*/ SHF.R.U32.HI R5, RZ, 0x1, R5; /* 0x00000001ff057819 */ - /* 0x000fe20000011605 */ - /*0330*/ IMAD R7, R7, 0x8, R4; /* 0x0000000807077824 */ - /* 0x000fe200078e0204 */ - /*0340*/ LOP3.LUT R6, R6, 0x1, RZ, 0xc0, !PT; /* 0x0000000106067812 */ - /* 0x000fc400078ec0ff */ - /*0350*/ LEA R4, R5, R4, 0x3; /* 0x0000000405047211 */ - /* 0x000fe200078e18ff */ - /*0360*/ IMAD.WIDE R12, R28, R10, c[0x0][0x168]; /* 0x00005a001c0c7625 */ - /* 0x000fe400078e020a */ - /*0370*/ IMAD R7, R6.reuse, 0x4, R7; /* 0x0000000406077824 */ - /* 0x040fe400078e0207 */ - /*0380*/ IMAD R4, R6, 0x4, R4; /* 0x0000000406047824 */ - /* 0x000fe400078e0204 */ - /*0390*/ IMAD.U32 R7, R7, 0x2, RZ; /* 0x0000000207077824 */ - /* 0x000fe400078e00ff */ - /*03a0*/ IMAD.U32 R5, R4, 0x2, RZ; /* 0x0000000204057824 */ - /* 0x000fe400078e00ff */ - /*03b0*/ IMAD.WIDE.U32 R16, R7, c[0x0][0x178], R16; /* 0x00005e0007107a25 */ - /* 0x000fc400078e0010 */ - /*03c0*/ IMAD.WIDE.U32 R12, R5, c[0x0][0x180], R12; /* 0x00006000050c7a25 */ - /* 0x000fd000078e000c */ - /*03d0*/ LD.E.128.SYS R24, [R16]; /* 0x0000000010187980 */ - /* 0x000064000010ed00 */ - /*03e0*/ LD.E.128.SYS R20, [R12]; /* 0x000000000c147980 */ - /* 0x000462000010ed00 */ - /*03f0*/ LD.E.128.SYS R16, [R16+0x10]; /* 0x0000001010107980 */ - /* 0x001e22000010ed00 */ - /*0400*/ LD.E.128.SYS R12, [R12+0x10]; /* 0x000000100c0c7980 */ - /* 0x004e22000010ed00 */ - /*0410*/ IMAD.U32 R8, RZ, RZ, RZ; /* 0x000000ffff087224 */ - /* 0x000fe200078e00ff */ - /*0420*/ MOV R9, RZ; /* 0x000000ff00097202 */ - /* 0x000fe20000000f00 */ - /*0430*/ IMAD.U32 R10, RZ, RZ, RZ; /* 0x000000ffff0a7224 */ - /* 0x000fe400078e00ff */ - /*0440*/ IMAD.U32 R11, RZ, RZ, RZ; /* 0x000000ffff0b7224 */ - /* 0x000fe200078e00ff */ - /*0450*/ MOV R7, RZ; /* 0x000000ff00077202 */ - /* 0x000fe20000000f00 */ - /*0460*/ IMAD.U32 R4, RZ, RZ, RZ; /* 0x000000ffff047224 */ - /* 0x000fc400078e00ff */ - /*0470*/ IMAD.U32 R5, RZ, RZ, RZ; /* 0x000000ffff057224 */ - /* 0x000fe400078e00ff */ - /*0480*/ IMAD.U32 R6, RZ, RZ, RZ; /* 0x000000ffff067224 */ - /* 0x000fe200078e00ff */ - /*0490*/ WARPSYNC 0xffffffff; /* 0xffffffff00007948 */ - /* 0x000fe20003800000 */ - /*04a0*/ HMMA.884.F32.F32.STEP0 R8, R24.reuse, R20.reuse.T, R8; /* 0x0000001418087236 */ - /* 0x0c22640000005408 */ - /*04b0*/ HMMA.884.F32.F32.STEP1 R10, R24.reuse, R20.reuse.T, R10; /* 0x00000014180a7236 */ - /* 0x0c04a4000000d40a */ - /*04c0*/ HMMA.884.F32.F32.STEP2 R4, R24.reuse, R20.reuse.T, R4; /* 0x0000001418047236 */ - /* 0x0c06e40000015404 */ - /*04d0*/ HMMA.884.F32.F32.STEP3 R6, R24, R20.T, R6; /* 0x0000001418067236 */ - /* 0x000928000001d406 */ - /*04e0*/ HMMA.884.F32.F32.STEP0 R8, R26.reuse, R22.reuse.T, R8; /* 0x000000161a087236 */ - /* 0x0c20240000005408 */ - /*04f0*/ HMMA.884.F32.F32.STEP1 R10, R26.reuse, R22.reuse.T, R10; /* 0x000000161a0a7236 */ - /* 0x0c4264000000d40a */ - /*0500*/ HMMA.884.F32.F32.STEP2 R4, R26.reuse, R22.reuse.T, R4; /* 0x000000161a047236 */ - /* 0x0c84a40000015404 */ - /*0510*/ HMMA.884.F32.F32.STEP3 R6, R26, R22.T, R6; /* 0x000000161a067236 */ - /* 0x0106e8000001d406 */ - /*0520*/ HMMA.884.F32.F32.STEP0 R8, R16.reuse, R12.reuse.T, R8; /* 0x0000000c10087236 */ - /* 0x0c10240000005408 */ - /*0530*/ HMMA.884.F32.F32.STEP1 R10, R16.reuse, R12.reuse.T, R10; /* 0x0000000c100a7236 */ - /* 0x0c2264000000d40a */ - /*0540*/ HMMA.884.F32.F32.STEP2 R4, R16.reuse, R12.reuse.T, R4; /* 0x0000000c10047236 */ - /* 0x0c44a40000015404 */ - /*0550*/ HMMA.884.F32.F32.STEP3 R6, R16, R12.T, R6; /* 0x0000000c10067236 */ - /* 0x0086e8000001d406 */ - /*0560*/ HMMA.884.F32.F32.STEP0 R8, R18.reuse, R14.reuse.T, R8; /* 0x0000000e12087236 */ - /* 0x0c10240000005408 */ - /*0570*/ HMMA.884.F32.F32.STEP1 R10, R18.reuse, R14.reuse.T, R10; /* 0x0000000e120a7236 */ - /* 0x0c2024000000d40a */ - /*0580*/ HMMA.884.F32.F32.STEP2 R4, R18.reuse, R14.reuse.T, R4; /* 0x0000000e12047236 */ - /* 0x0c40240000015404 */ - /*0590*/ HMMA.884.F32.F32.STEP3 R6, R18, R14.T, R6; /* 0x0000000e12067236 */ - /* 0x008034000001d406 */ - /*05a0*/ BSYNC B0; /* 0x0000000000007941 */ - /* 0x001fea0003800000 */ - /*05b0*/ S2R R12, SR_LANEID; /* 0x00000000000c7919 */ - /* 0x000e220000000000 */ - /*05c0*/ IMAD R28, R28, c[0x0][0x178], RZ; /* 0x00005e001c1c7a24 */ - /* 0x000fe200078e02ff */ - /*05d0*/ SHF.R.U32.HI R14, RZ, 0x4, R12; /* 0x00000004ff0e7819 */ - /* 0x001fc4000001160c */ - /*05e0*/ SHF.R.U32.HI R13, RZ, 0x2, R12; /* 0x00000002ff0d7819 */ - /* 0x000fe4000001160c */ - /*05f0*/ LOP3.LUT R12, R12, 0x3, RZ, 0xc0, !PT; /* 0x000000030c0c7812 */ - /* 0x000fe400078ec0ff */ - /*0600*/ LOP3.LUT R14, R14, 0x1, RZ, 0xc0, !PT; /* 0x000000010e0e7812 */ - /* 0x000fe400078ec0ff */ - /*0610*/ LOP3.LUT R13, R13, 0x3, RZ, 0xc0, !PT; /* 0x000000030d0d7812 */ - /* 0x000fc600078ec0ff */ - /*0620*/ IMAD R12, R14, 0x4, R12; /* 0x000000040e0c7824 */ - /* 0x000fe200078e020c */ - /*0630*/ LOP3.LUT R15, R13, 0x1, RZ, 0xc0, !PT; /* 0x000000010d0f7812 */ - /* 0x000fe400078ec0ff */ - /*0640*/ SHF.R.U32.HI R16, RZ, 0x1, R13; /* 0x00000001ff107819 */ - /* 0x000fe4000001160d */ - /*0650*/ LOP3.LUT R13, R12.reuse, 0x5, RZ, 0xc0, !PT; /* 0x000000050c0d7812 */ - /* 0x040fe400078ec0ff */ - /*0660*/ LOP3.LUT R14, R12, 0x2, RZ, 0xc0, !PT; /* 0x000000020c0e7812 */ - /* 0x000fc600078ec0ff */ - /*0670*/ IMAD R12, R15, 0x8, R13; /* 0x000000080f0c7824 */ - /* 0x000fe200078e020d */ - /*0680*/ SHF.R.S32.HI R15, RZ, 0x1f, R29; /* 0x0000001fff0f7819 */ - /* 0x000fe2000001141d */ - /*0690*/ IMAD R14, R16, 0x8, R14; /* 0x00000008100e7824 */ - /* 0x000fe200078e020e */ - /*06a0*/ IADD3 R17, P0, R29, R28, RZ; /* 0x0000001c1d117210 */ - /* 0x000fe20007f1e0ff */ - /*06b0*/ IMAD.U32 R13, RZ, RZ, RZ; /* 0x000000ffff0d7224 */ - /* 0x000fc600078e00ff */ - /*06c0*/ LEA.HI.X.SX32 R15, R28, R15, 0x1, P0; /* 0x0000000f1c0f7211 */ - /* 0x000fe200000f0eff */ - /*06d0*/ IMAD.WIDE.U32 R12, R14, c[0x0][0x178], R12; /* 0x00005e000e0c7a25 */ - /* 0x000fe200078e000c */ - /*06e0*/ LEA R14, P0, R17, c[0x0][0x170], 0x2; /* 0x00005c00110e7a11 */ - /* 0x000fe200078010ff */ - /*06f0*/ IMAD.U32 R18, RZ, RZ, c[0x0][0x178]; /* 0x00005e00ff127624 */ - /* 0x000fc600078e00ff */ - /*0700*/ LEA.HI.X R17, R17, c[0x0][0x174], R15, 0x2, P0; /* 0x00005d0011117a11 */ - /* 0x000fe400000f140f */ - /*0710*/ LEA R16, P0, R12, R14, 0x2; /* 0x0000000e0c107211 */ - /* 0x000fe400078010ff */ - /*0720*/ SHF.L.U32 R14, R18, 0x2, RZ; /* 0x00000002120e7819 */ - /* 0x000fe400000006ff */ - /*0730*/ SHF.R.U32.HI R15, RZ, 0x1e, R18; /* 0x0000001eff0f7819 */ - /* 0x000fe40000011612 */ - /*0740*/ LEA.HI.X R13, R12, R17, R13, 0x2, P0; /* 0x000000110c0d7211 */ - /* 0x000fe400000f140d */ - /*0750*/ LEA R18, P1, R14, R16, 0x2; /* 0x000000100e127211 */ - /* 0x000fc400078210ff */ - /*0760*/ IADD3 R17, P0, R14.reuse, R16, RZ; /* 0x000000100e117210 */ - /* 0x040fe40007f1e0ff */ - /*0770*/ LEA.HI.X R19, R14.reuse, R13, R15, 0x2, P1; /* 0x0000000d0e137211 */ - /* 0x040fe400008f140f */ - /*0780*/ IADD3 R21, P1, R14, R18, RZ; /* 0x000000120e157210 */ - /* 0x000fe20007f3e0ff */ - /*0790*/ IMAD.X R20, R15.reuse, 0x1, R13, P0; /* 0x000000010f147824 */ - /* 0x040fe400000e060d */ - /*07a0*/ IMAD.U32 R12, RZ, RZ, R16; /* 0x000000ffff0c7224 */ - /* 0x000fe400078e0010 */ - /*07b0*/ IMAD.X R22, R15, 0x1, R19, P1; /* 0x000000010f167824 */ - /* 0x000fe400008e0613 */ - /*07c0*/ IMAD.U32 R14, RZ, RZ, R17; /* 0x000000ffff0e7224 */ - /* 0x000fc400078e0011 */ - /*07d0*/ IMAD.U32 R15, RZ, RZ, R20; /* 0x000000ffff0f7224 */ - /* 0x000fe200078e0014 */ - /*07e0*/ MOV R16, R21; /* 0x0000001500107202 */ - /* 0x000fe20000000f00 */ - /*07f0*/ IMAD.U32 R17, RZ, RZ, R22; /* 0x000000ffff117224 */ - /* 0x000fe200078e0016 */ - /*0800*/ ST.E.SYS [R12], R8; /* 0x000000000c007385 */ - /* 0x0001e2000010e908 */ - /*0810*/ ST.E.SYS [R12+0x8], R10; /* 0x000000080c007385 */ - /* 0x0003e8000010e90a */ - /*0820*/ ST.E.SYS [R14], R9; /* 0x000000000e007385 */ - /* 0x0003e2000010e909 */ - /*0830*/ ST.E.SYS [R14+0x8], R11; /* 0x000000080e007385 */ - /* 0x0003e2000010e90b */ - /*0840*/ ST.E.SYS [R18], R4; /* 0x0000000012007385 */ - /* 0x0003e2000010e904 */ - /*0850*/ ST.E.SYS [R18+0x8], R6; /* 0x0000000812007385 */ - /* 0x0003e2000010e906 */ - /*0860*/ ST.E.SYS [R16], R5; /* 0x0000000010007385 */ - /* 0x0003e2000010e905 */ - /*0870*/ ST.E.SYS [R16+0x8], R7; /* 0x0000000810007385 */ - /* 0x0003e2000010e907 */ - /*0880*/ WARPSYNC 0xffffffff; /* 0xffffffff00007948 */ - /* 0x000fe20003800000 */ - /*0890*/ IADD3 R8, R2, -c[0x0][0x20], RZ; /* 0x8000080002087a10 */ - /* 0x001fd00007ffe0ff */ - /*08a0*/ CS2R.32 R4, SR_CLOCKLO; /* 0x0000000000047805 */ - /* 0x002fd00000005000 */ - /*08b0*/ IMAD R3, R4, 0x1, -R3; /* 0x0000000104037824 */ - /* 0x000fd000078e0a03 */ - /*08c0*/ STL [R8], R3; /* 0x0000000308007387 */ - /* 0x0001e20000100800 */ - /*08d0*/ IMAD.U32 R6, RZ, RZ, R2; /* 0x000000ffff067224 */ - /* 0x000fe200078e0002 */ - /*08e0*/ MOV R4, 0x0; /* 0x0000000000047802 */ - /* 0x000fe20000000f00 */ - /*08f0*/ IMAD.U32 R7, RZ, RZ, R0; /* 0x000000ffff077224 */ - /* 0x000fe200078e0000 */ - /*0900*/ MOV R5, 0x0; /* 0x0000000000057802 */ - /* 0x000fe40000000f00 */ - /*0910*/ MOV R20, 0x0; /* 0x0000000000147802 */ - /* 0x000fe40000000f00 */ - /*0920*/ MOV R21, 0x0; /* 0x0000000000157802 */ - /* 0x000fd00000000f00 */ - /*0930*/ CALL.ABS.NOINC 0x0; /* 0x0000000000007943 */ - /* 0x001fea0003c00000 */ - /*0940*/ EXIT; /* 0x000000000000794d */ - /* 0x000fea0003800000 */ - /*0950*/ BRA 0x950; /* 0xfffffff000007947 */ - /* 0x000fc0000383ffff */ - /*0960*/ NOP; /* 0x0000000000007918 */ - /* 0x000fc00000000000 */ - /*0970*/ NOP; /* 0x0000000000007918 */ - /* 0x000fc00000000000 */ - ............................................. - - - -Fatbin ptx code: -================ -arch = sm_70 -code version = [6,0] -producer = cuda -host = linux -compile_size = 64bit -compressed - - - - - - - - -.version 6.0 -.target sm_70 -.address_size 64 - - -.extern .func (.param .b32 func_retval0) vprintf -( -.param .b64 vprintf_param_0, -.param .b64 vprintf_param_1 -) -; -.global .align 16 .b8 $str[9] = {99, 108, 111, 99, 107, 61, 37, 100, 0}; - -.visible .entry _Z12wmma_exampleP6__halfS0_Pfiiiff( -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_0, -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_1, -.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_2, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_3, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_4, -.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_5, -.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_6, -.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_7 -) -{ -.local .align 8 .b8 __local_depot0[8]; -.reg .b64 %SP; -.reg .b64 %SPL; -.reg .pred %p<6>; -.reg .f32 %f<34>; -.reg .b32 %r<38>; -.reg .b64 %rd<18>; - - -mov.u64 %rd17, __local_depot0; -cvta.local.u64 %SP, %rd17; -ld.param.u64 %rd1, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0]; -ld.param.u64 %rd2, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1]; -ld.param.u64 %rd3, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2]; -ld.param.u32 %r4, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3]; -ld.param.u32 %r7, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4]; -ld.param.u32 %r5, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5]; - - mov.u32 %r6, %clock; - - mov.u32 %r8, %ntid.x; -mov.u32 %r9, %ctaid.x; -mov.u32 %r10, %tid.x; -mad.lo.s32 %r11, %r8, %r9, %r10; -mov.u32 %r12, WARP_SZ; -div.u32 %r13, %r11, %r12; -mov.u32 %r14, %ntid.y; -mov.u32 %r15, %ctaid.y; -mov.u32 %r16, %tid.y; -mad.lo.s32 %r17, %r14, %r15, %r16; -shl.b32 %r2, %r13, 4; -shl.b32 %r3, %r17, 4; -setp.lt.s32 %p1, %r2, %r4; -setp.gt.s32 %p2, %r5, 0; -and.pred %p3, %p1, %p2; -setp.lt.s32 %p4, %r3, %r7; -and.pred %p5, %p3, %p4; -mov.f32 %f26, 0f00000000; -mov.f32 %f27, %f26; -mov.f32 %f28, %f26; -mov.f32 %f29, %f26; -mov.f32 %f30, %f26; -mov.f32 %f31, %f26; -mov.f32 %f32, %f26; -mov.f32 %f33, %f26; -@!%p5 bra BB0_2; -bra.uni BB0_1; - -BB0_1: -mul.wide.s32 %rd4, %r2, 2; -add.s64 %rd5, %rd1, %rd4; -wmma.load.a.sync.row.m16n16k16.f16 {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, [%rd5], %r4; -mul.wide.s32 %rd6, %r3, 2; -add.s64 %rd7, %rd2, %rd6; -wmma.load.b.sync.col.m16n16k16.f16 {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, [%rd7], %r5; -mov.f32 %f25, 0f00000000; -wmma.mma.sync.row.col.m16n16k16.f32.f32 {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, {%f25, %f25, %f25, %f25, %f25, %f25, %f25, %f25}; - -BB0_2: -add.u64 %rd8, %SP, 0; -cvta.to.local.u64 %rd9, %rd8; -mul.lo.s32 %r35, %r3, %r4; -cvt.s64.s32 %rd10, %r35; -cvt.s64.s32 %rd11, %r2; -add.s64 %rd12, %rd10, %rd11; -shl.b64 %rd13, %rd12, 2; -add.s64 %rd14, %rd3, %rd13; -wmma.store.d.sync.col.m16n16k16.f32 [%rd14], {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, %r4; - - mov.u32 %r34, %clock; - - sub.s32 %r36, %r34, %r6; -st.local.u32 [%rd9], %r36; -mov.u64 %rd15, $str; -cvta.global.u64 %rd16, %rd15; - - { -.reg .b32 temp_param_reg; - - .param .b64 param0; -st.param.b64 [param0+0], %rd16; -.param .b64 param1; -st.param.b64 [param1+0], %rd8; -.param .b32 retval0; -call.uni (retval0), -vprintf, -( -param0, -param1 -); -ld.param.b32 %r37, [retval0+0]; - - - } - ret; -} - - -.visible .entry _Z17convertFp32ToFp16P6__halfPfi( -.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_0, -.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_1, -.param .u32 _Z17convertFp32ToFp16P6__halfPfi_param_2 -) -{ -.reg .pred %p<2>; -.reg .b16 %rs<2>; -.reg .f32 %f<2>; -.reg .b32 %r<6>; -.reg .b64 %rd<9>; - - -ld.param.u64 %rd1, [_Z17convertFp32ToFp16P6__halfPfi_param_0]; -ld.param.u64 %rd2, [_Z17convertFp32ToFp16P6__halfPfi_param_1]; -ld.param.u32 %r2, [_Z17convertFp32ToFp16P6__halfPfi_param_2]; -mov.u32 %r3, %ntid.x; -mov.u32 %r4, %ctaid.x; -mov.u32 %r5, %tid.x; -mad.lo.s32 %r1, %r4, %r3, %r5; -setp.ge.s32 %p1, %r1, %r2; -@%p1 bra BB1_2; - -cvta.to.global.u64 %rd3, %rd2; -mul.wide.s32 %rd4, %r1, 4; -add.s64 %rd5, %rd3, %rd4; -ld.global.f32 %f1, [%rd5]; - - { cvt.rn.f16.f32 %rs1, %f1;} - - - cvta.to.global.u64 %rd6, %rd1; -mul.wide.s32 %rd7, %r1, 2; -add.s64 %rd8, %rd6, %rd7; -st.global.u16 [%rd8], %rs1; - -BB1_2: -ret; -} - - diff --git a/cuda-kernels/gpgpu_inst_stats.txt b/cuda-kernels/gpgpu_inst_stats.txt deleted file mode 100755 index 41f06a4..0000000 --- a/cuda-kernels/gpgpu_inst_stats.txt +++ /dev/null @@ -1,20 +0,0 @@ -kernel line : count latency dram_traffic smem_bk_conflicts smem_warp gmem_access_generated gmem_warp exposed_latency warp_divergence -_1.ptx 164 : 512 2560 1024 0 0 16 16 0 0 -_1.ptx 163 : 512 5696 0 0 0 0 0 0 0 -_1.ptx 162 : 512 5664 0 0 0 0 0 0 0 -_1.ptx 161 : 512 3616 0 0 0 0 0 0 0 -_1.ptx 158 : 512 3616 0 0 0 0 0 0 0 -_1.ptx 156 : 512 134048 2048 0 0 16 16 0 0 -_1.ptx 155 : 512 5632 0 0 0 0 0 0 0 -_1.ptx 154 : 512 5376 0 0 0 0 0 0 0 -_1.ptx 143 : 512 100864 128 0 0 0 0 0 0 -_1.ptx 167 : 512 3072 0 0 0 0 0 0 0 -_1.ptx 144 : 512 100864 0 0 0 0 0 0 0 -_1.ptx 145 : 512 1536 0 0 0 0 0 0 0 -_1.ptx 146 : 512 3072 0 0 0 0 0 0 0 -_1.ptx 147 : 512 3072 0 0 0 0 0 0 0 -_1.ptx 148 : 512 3072 0 0 0 0 0 0 0 -_1.ptx 149 : 512 8384 0 0 0 0 0 0 0 -_1.ptx 150 : 512 4096 0 0 0 0 0 0 0 -_1.ptx 151 : 512 0 0 0 0 0 0 0 0 -_1.ptx 153 : 512 3968 0 0 0 0 0 0 0 diff --git a/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-34-2018.log b/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-34-2018.log deleted file mode 100644 index f754f0c..0000000 --- a/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-34-2018.log +++ /dev/null @@ -1,324 +0,0 @@ -kernel_name = -kernel_launch_uid = - -Kernel Average Power Data: -kernel_avg_power = 0 -gpu_avg_IBP, = -nan -gpu_avg_ICP, = -nan -gpu_avg_DCP, = -nan -gpu_avg_TCP, = -nan -gpu_avg_CCP, = -nan -gpu_avg_SHRDP, = -nan -gpu_avg_RFP, = -nan -gpu_avg_SPP, = -nan -gpu_avg_SFUP, = -nan -gpu_avg_FPUP, = -nan -gpu_avg_SCHEDP, = -nan -gpu_avg_L2CP, = -nan -gpu_avg_MCP, = -nan -gpu_avg_NOCP, = -nan -gpu_avg_DRAMP, = -nan -gpu_avg_PIPEP, = -nan -gpu_avg_IDLE_COREP, = -nan -gpu_avg_CONST_DYNAMICP = -nan -gpu_avg_TOT_INST, = -nan -gpu_avg_FP_INT, = -nan -gpu_avg_IC_H, = -nan -gpu_avg_IC_M, = -nan -gpu_avg_DC_RH, = -nan -gpu_avg_DC_RM, = -nan -gpu_avg_DC_WH, = -nan -gpu_avg_DC_WM, = -nan -gpu_avg_TC_H, = -nan -gpu_avg_TC_M, = -nan -gpu_avg_CC_H, = -nan -gpu_avg_CC_M, = -nan -gpu_avg_SHRD_ACC, = -nan -gpu_avg_REG_RD, = -nan -gpu_avg_REG_WR, = -nan -gpu_avg_NON_REG_OPs, = -nan -gpu_avg_SP_ACC, = -nan -gpu_avg_SFU_ACC, = -nan -gpu_avg_FPU_ACC, = -nan -gpu_avg_MEM_RD, = -nan -gpu_avg_MEM_WR, = -nan -gpu_avg_MEM_PRE, = -nan -gpu_avg_L2_RH, = -nan -gpu_avg_L2_RM, = -nan -gpu_avg_L2_WH, = -nan -gpu_avg_L2_WM, = -nan -gpu_avg_NOC_A, = -nan -gpu_avg_PIPE_A, = -nan -gpu_avg_IDLE_CORE_N, = -nan -gpu_avg_CONST_DYNAMICN = -nan - -Kernel Maximum Power Data: -kernel_max_power = 0 -gpu_max_IBP, = 0 -gpu_max_ICP, = 0 -gpu_max_DCP, = 0 -gpu_max_TCP, = 0 -gpu_max_CCP, = 0 -gpu_max_SHRDP, = 0 -gpu_max_RFP, = 0 -gpu_max_SPP, = 0 -gpu_max_SFUP, = 0 -gpu_max_FPUP, = 0 -gpu_max_SCHEDP, = 0 -gpu_max_L2CP, = 0 -gpu_max_MCP, = 0 -gpu_max_NOCP, = 0 -gpu_max_DRAMP, = 0 -gpu_max_PIPEP, = 0 -gpu_max_IDLE_COREP, = 0 -gpu_max_CONST_DYNAMICP = 0 -gpu_max_TOT_INST, = 0 -gpu_max_FP_INT, = 0 -gpu_max_IC_H, = 0 -gpu_max_IC_M, = 0 -gpu_max_DC_RH, = 0 -gpu_max_DC_RM, = 0 -gpu_max_DC_WH, = 0 -gpu_max_DC_WM, = 0 -gpu_max_TC_H, = 0 -gpu_max_TC_M, = 0 -gpu_max_CC_H, = 0 -gpu_max_CC_M, = 0 -gpu_max_SHRD_ACC, = 0 -gpu_max_REG_RD, = 0 -gpu_max_REG_WR, = 0 -gpu_max_NON_REG_OPs, = 0 -gpu_max_SP_ACC, = 0 -gpu_max_SFU_ACC, = 0 -gpu_max_FPU_ACC, = 0 -gpu_max_MEM_RD, = 0 -gpu_max_MEM_WR, = 0 -gpu_max_MEM_PRE, = 0 -gpu_max_L2_RH, = 0 -gpu_max_L2_RM, = 0 -gpu_max_L2_WH, = 0 -gpu_max_L2_WM, = 0 -gpu_max_NOC_A, = 0 -gpu_max_PIPE_A, = 0 -gpu_max_IDLE_CORE_N, = 0 -gpu_max_CONST_DYNAMICN = 0 - -Kernel Minimum Power Data: -kernel_min_power = 0 -gpu_min_IBP, = 0 -gpu_min_ICP, = 0 -gpu_min_DCP, = 0 -gpu_min_TCP, = 0 -gpu_min_CCP, = 0 -gpu_min_SHRDP, = 0 -gpu_min_RFP, = 0 -gpu_min_SPP, = 0 -gpu_min_SFUP, = 0 -gpu_min_FPUP, = 0 -gpu_min_SCHEDP, = 0 -gpu_min_L2CP, = 0 -gpu_min_MCP, = 0 -gpu_min_NOCP, = 0 -gpu_min_DRAMP, = 0 -gpu_min_PIPEP, = 0 -gpu_min_IDLE_COREP, = 0 -gpu_min_CONST_DYNAMICP = 0 -gpu_min_TOT_INST, = 0 -gpu_min_FP_INT, = 0 -gpu_min_IC_H, = 0 -gpu_min_IC_M, = 0 -gpu_min_DC_RH, = 0 -gpu_min_DC_RM, = 0 -gpu_min_DC_WH, = 0 -gpu_min_DC_WM, = 0 -gpu_min_TC_H, = 0 -gpu_min_TC_M, = 0 -gpu_min_CC_H, = 0 -gpu_min_CC_M, = 0 -gpu_min_SHRD_ACC, = 0 -gpu_min_REG_RD, = 0 -gpu_min_REG_WR, = 0 -gpu_min_NON_REG_OPs, = 0 -gpu_min_SP_ACC, = 0 -gpu_min_SFU_ACC, = 0 -gpu_min_FPU_ACC, = 0 -gpu_min_MEM_RD, = 0 -gpu_min_MEM_WR, = 0 -gpu_min_MEM_PRE, = 0 -gpu_min_L2_RH, = 0 -gpu_min_L2_RM, = 0 -gpu_min_L2_WH, = 0 -gpu_min_L2_WM, = 0 -gpu_min_NOC_A, = 0 -gpu_min_PIPE_A, = 0 -gpu_min_IDLE_CORE_N, = 0 -gpu_min_CONST_DYNAMICN = 0 - -Accumulative Power Statistics Over Previous Kernels: -gpu_tot_avg_power = -nan -gpu_tot_max_power = 0 -gpu_tot_min_power = 0 - - -kernel_name = -kernel_launch_uid = - -Kernel Average Power Data: -kernel_avg_power = 0 -gpu_avg_IBP, = -nan -gpu_avg_ICP, = -nan -gpu_avg_DCP, = -nan -gpu_avg_TCP, = -nan -gpu_avg_CCP, = -nan -gpu_avg_SHRDP, = -nan -gpu_avg_RFP, = -nan -gpu_avg_SPP, = -nan -gpu_avg_SFUP, = -nan -gpu_avg_FPUP, = -nan -gpu_avg_SCHEDP, = -nan -gpu_avg_L2CP, = -nan -gpu_avg_MCP, = -nan -gpu_avg_NOCP, = -nan -gpu_avg_DRAMP, = -nan -gpu_avg_PIPEP, = -nan -gpu_avg_IDLE_COREP, = -nan -gpu_avg_CONST_DYNAMICP = -nan -gpu_avg_TOT_INST, = -nan -gpu_avg_FP_INT, = -nan -gpu_avg_IC_H, = -nan -gpu_avg_IC_M, = -nan -gpu_avg_DC_RH, = -nan -gpu_avg_DC_RM, = -nan -gpu_avg_DC_WH, = -nan -gpu_avg_DC_WM, = -nan -gpu_avg_TC_H, = -nan -gpu_avg_TC_M, = -nan -gpu_avg_CC_H, = -nan -gpu_avg_CC_M, = -nan -gpu_avg_SHRD_ACC, = -nan -gpu_avg_REG_RD, = -nan -gpu_avg_REG_WR, = -nan -gpu_avg_NON_REG_OPs, = -nan -gpu_avg_SP_ACC, = -nan -gpu_avg_SFU_ACC, = -nan -gpu_avg_FPU_ACC, = -nan -gpu_avg_MEM_RD, = -nan -gpu_avg_MEM_WR, = -nan -gpu_avg_MEM_PRE, = -nan -gpu_avg_L2_RH, = -nan -gpu_avg_L2_RM, = -nan -gpu_avg_L2_WH, = -nan -gpu_avg_L2_WM, = -nan -gpu_avg_NOC_A, = -nan -gpu_avg_PIPE_A, = -nan -gpu_avg_IDLE_CORE_N, = -nan -gpu_avg_CONST_DYNAMICN = -nan - -Kernel Maximum Power Data: -kernel_max_power = 0 -gpu_max_IBP, = 0 -gpu_max_ICP, = 0 -gpu_max_DCP, = 0 -gpu_max_TCP, = 0 -gpu_max_CCP, = 0 -gpu_max_SHRDP, = 0 -gpu_max_RFP, = 0 -gpu_max_SPP, = 0 -gpu_max_SFUP, = 0 -gpu_max_FPUP, = 0 -gpu_max_SCHEDP, = 0 -gpu_max_L2CP, = 0 -gpu_max_MCP, = 0 -gpu_max_NOCP, = 0 -gpu_max_DRAMP, = 0 -gpu_max_PIPEP, = 0 -gpu_max_IDLE_COREP, = 0 -gpu_max_CONST_DYNAMICP = 0 -gpu_max_TOT_INST, = 0 -gpu_max_FP_INT, = 0 -gpu_max_IC_H, = 0 -gpu_max_IC_M, = 0 -gpu_max_DC_RH, = 0 -gpu_max_DC_RM, = 0 -gpu_max_DC_WH, = 0 -gpu_max_DC_WM, = 0 -gpu_max_TC_H, = 0 -gpu_max_TC_M, = 0 -gpu_max_CC_H, = 0 -gpu_max_CC_M, = 0 -gpu_max_SHRD_ACC, = 0 -gpu_max_REG_RD, = 0 -gpu_max_REG_WR, = 0 -gpu_max_NON_REG_OPs, = 0 -gpu_max_SP_ACC, = 0 -gpu_max_SFU_ACC, = 0 -gpu_max_FPU_ACC, = 0 -gpu_max_MEM_RD, = 0 -gpu_max_MEM_WR, = 0 -gpu_max_MEM_PRE, = 0 -gpu_max_L2_RH, = 0 -gpu_max_L2_RM, = 0 -gpu_max_L2_WH, = 0 -gpu_max_L2_WM, = 0 -gpu_max_NOC_A, = 0 -gpu_max_PIPE_A, = 0 -gpu_max_IDLE_CORE_N, = 0 -gpu_max_CONST_DYNAMICN = 0 - -Kernel Minimum Power Data: -kernel_min_power = 0 -gpu_min_IBP, = 0 -gpu_min_ICP, = 0 -gpu_min_DCP, = 0 -gpu_min_TCP, = 0 -gpu_min_CCP, = 0 -gpu_min_SHRDP, = 0 -gpu_min_RFP, = 0 -gpu_min_SPP, = 0 -gpu_min_SFUP, = 0 -gpu_min_FPUP, = 0 -gpu_min_SCHEDP, = 0 -gpu_min_L2CP, = 0 -gpu_min_MCP, = 0 -gpu_min_NOCP, = 0 -gpu_min_DRAMP, = 0 -gpu_min_PIPEP, = 0 -gpu_min_IDLE_COREP, = 0 -gpu_min_CONST_DYNAMICP = 0 -gpu_min_TOT_INST, = 0 -gpu_min_FP_INT, = 0 -gpu_min_IC_H, = 0 -gpu_min_IC_M, = 0 -gpu_min_DC_RH, = 0 -gpu_min_DC_RM, = 0 -gpu_min_DC_WH, = 0 -gpu_min_DC_WM, = 0 -gpu_min_TC_H, = 0 -gpu_min_TC_M, = 0 -gpu_min_CC_H, = 0 -gpu_min_CC_M, = 0 -gpu_min_SHRD_ACC, = 0 -gpu_min_REG_RD, = 0 -gpu_min_REG_WR, = 0 -gpu_min_NON_REG_OPs, = 0 -gpu_min_SP_ACC, = 0 -gpu_min_SFU_ACC, = 0 -gpu_min_FPU_ACC, = 0 -gpu_min_MEM_RD, = 0 -gpu_min_MEM_WR, = 0 -gpu_min_MEM_PRE, = 0 -gpu_min_L2_RH, = 0 -gpu_min_L2_RM, = 0 -gpu_min_L2_WH, = 0 -gpu_min_L2_WM, = 0 -gpu_min_NOC_A, = 0 -gpu_min_PIPE_A, = 0 -gpu_min_IDLE_CORE_N, = 0 -gpu_min_CONST_DYNAMICN = 0 - -Accumulative Power Statistics Over Previous Kernels: -gpu_tot_avg_power = -nan -gpu_tot_max_power = 0 -gpu_tot_min_power = 0 - - diff --git a/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-47-2018.log b/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-47-2018.log deleted file mode 100644 index f754f0c..0000000 --- a/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-47-2018.log +++ /dev/null @@ -1,324 +0,0 @@ -kernel_name = -kernel_launch_uid = - -Kernel Average Power Data: -kernel_avg_power = 0 -gpu_avg_IBP, = -nan -gpu_avg_ICP, = -nan -gpu_avg_DCP, = -nan -gpu_avg_TCP, = -nan -gpu_avg_CCP, = -nan -gpu_avg_SHRDP, = -nan -gpu_avg_RFP, = -nan -gpu_avg_SPP, = -nan -gpu_avg_SFUP, = -nan -gpu_avg_FPUP, = -nan -gpu_avg_SCHEDP, = -nan -gpu_avg_L2CP, = -nan -gpu_avg_MCP, = -nan -gpu_avg_NOCP, = -nan -gpu_avg_DRAMP, = -nan -gpu_avg_PIPEP, = -nan -gpu_avg_IDLE_COREP, = -nan -gpu_avg_CONST_DYNAMICP = -nan -gpu_avg_TOT_INST, = -nan -gpu_avg_FP_INT, = -nan -gpu_avg_IC_H, = -nan -gpu_avg_IC_M, = -nan -gpu_avg_DC_RH, = -nan -gpu_avg_DC_RM, = -nan -gpu_avg_DC_WH, = -nan -gpu_avg_DC_WM, = -nan -gpu_avg_TC_H, = -nan -gpu_avg_TC_M, = -nan -gpu_avg_CC_H, = -nan -gpu_avg_CC_M, = -nan -gpu_avg_SHRD_ACC, = -nan -gpu_avg_REG_RD, = -nan -gpu_avg_REG_WR, = -nan -gpu_avg_NON_REG_OPs, = -nan -gpu_avg_SP_ACC, = -nan -gpu_avg_SFU_ACC, = -nan -gpu_avg_FPU_ACC, = -nan -gpu_avg_MEM_RD, = -nan -gpu_avg_MEM_WR, = -nan -gpu_avg_MEM_PRE, = -nan -gpu_avg_L2_RH, = -nan -gpu_avg_L2_RM, = -nan -gpu_avg_L2_WH, = -nan -gpu_avg_L2_WM, = -nan -gpu_avg_NOC_A, = -nan -gpu_avg_PIPE_A, = -nan -gpu_avg_IDLE_CORE_N, = -nan -gpu_avg_CONST_DYNAMICN = -nan - -Kernel Maximum Power Data: -kernel_max_power = 0 -gpu_max_IBP, = 0 -gpu_max_ICP, = 0 -gpu_max_DCP, = 0 -gpu_max_TCP, = 0 -gpu_max_CCP, = 0 -gpu_max_SHRDP, = 0 -gpu_max_RFP, = 0 -gpu_max_SPP, = 0 -gpu_max_SFUP, = 0 -gpu_max_FPUP, = 0 -gpu_max_SCHEDP, = 0 -gpu_max_L2CP, = 0 -gpu_max_MCP, = 0 -gpu_max_NOCP, = 0 -gpu_max_DRAMP, = 0 -gpu_max_PIPEP, = 0 -gpu_max_IDLE_COREP, = 0 -gpu_max_CONST_DYNAMICP = 0 -gpu_max_TOT_INST, = 0 -gpu_max_FP_INT, = 0 -gpu_max_IC_H, = 0 -gpu_max_IC_M, = 0 -gpu_max_DC_RH, = 0 -gpu_max_DC_RM, = 0 -gpu_max_DC_WH, = 0 -gpu_max_DC_WM, = 0 -gpu_max_TC_H, = 0 -gpu_max_TC_M, = 0 -gpu_max_CC_H, = 0 -gpu_max_CC_M, = 0 -gpu_max_SHRD_ACC, = 0 -gpu_max_REG_RD, = 0 -gpu_max_REG_WR, = 0 -gpu_max_NON_REG_OPs, = 0 -gpu_max_SP_ACC, = 0 -gpu_max_SFU_ACC, = 0 -gpu_max_FPU_ACC, = 0 -gpu_max_MEM_RD, = 0 -gpu_max_MEM_WR, = 0 -gpu_max_MEM_PRE, = 0 -gpu_max_L2_RH, = 0 -gpu_max_L2_RM, = 0 -gpu_max_L2_WH, = 0 -gpu_max_L2_WM, = 0 -gpu_max_NOC_A, = 0 -gpu_max_PIPE_A, = 0 -gpu_max_IDLE_CORE_N, = 0 -gpu_max_CONST_DYNAMICN = 0 - -Kernel Minimum Power Data: -kernel_min_power = 0 -gpu_min_IBP, = 0 -gpu_min_ICP, = 0 -gpu_min_DCP, = 0 -gpu_min_TCP, = 0 -gpu_min_CCP, = 0 -gpu_min_SHRDP, = 0 -gpu_min_RFP, = 0 -gpu_min_SPP, = 0 -gpu_min_SFUP, = 0 -gpu_min_FPUP, = 0 -gpu_min_SCHEDP, = 0 -gpu_min_L2CP, = 0 -gpu_min_MCP, = 0 -gpu_min_NOCP, = 0 -gpu_min_DRAMP, = 0 -gpu_min_PIPEP, = 0 -gpu_min_IDLE_COREP, = 0 -gpu_min_CONST_DYNAMICP = 0 -gpu_min_TOT_INST, = 0 -gpu_min_FP_INT, = 0 -gpu_min_IC_H, = 0 -gpu_min_IC_M, = 0 -gpu_min_DC_RH, = 0 -gpu_min_DC_RM, = 0 -gpu_min_DC_WH, = 0 -gpu_min_DC_WM, = 0 -gpu_min_TC_H, = 0 -gpu_min_TC_M, = 0 -gpu_min_CC_H, = 0 -gpu_min_CC_M, = 0 -gpu_min_SHRD_ACC, = 0 -gpu_min_REG_RD, = 0 -gpu_min_REG_WR, = 0 -gpu_min_NON_REG_OPs, = 0 -gpu_min_SP_ACC, = 0 -gpu_min_SFU_ACC, = 0 -gpu_min_FPU_ACC, = 0 -gpu_min_MEM_RD, = 0 -gpu_min_MEM_WR, = 0 -gpu_min_MEM_PRE, = 0 -gpu_min_L2_RH, = 0 -gpu_min_L2_RM, = 0 -gpu_min_L2_WH, = 0 -gpu_min_L2_WM, = 0 -gpu_min_NOC_A, = 0 -gpu_min_PIPE_A, = 0 -gpu_min_IDLE_CORE_N, = 0 -gpu_min_CONST_DYNAMICN = 0 - -Accumulative Power Statistics Over Previous Kernels: -gpu_tot_avg_power = -nan -gpu_tot_max_power = 0 -gpu_tot_min_power = 0 - - -kernel_name = -kernel_launch_uid = - -Kernel Average Power Data: -kernel_avg_power = 0 -gpu_avg_IBP, = -nan -gpu_avg_ICP, = -nan -gpu_avg_DCP, = -nan -gpu_avg_TCP, = -nan -gpu_avg_CCP, = -nan -gpu_avg_SHRDP, = -nan -gpu_avg_RFP, = -nan -gpu_avg_SPP, = -nan -gpu_avg_SFUP, = -nan -gpu_avg_FPUP, = -nan -gpu_avg_SCHEDP, = -nan -gpu_avg_L2CP, = -nan -gpu_avg_MCP, = -nan -gpu_avg_NOCP, = -nan -gpu_avg_DRAMP, = -nan -gpu_avg_PIPEP, = -nan -gpu_avg_IDLE_COREP, = -nan -gpu_avg_CONST_DYNAMICP = -nan -gpu_avg_TOT_INST, = -nan -gpu_avg_FP_INT, = -nan -gpu_avg_IC_H, = -nan -gpu_avg_IC_M, = -nan -gpu_avg_DC_RH, = -nan -gpu_avg_DC_RM, = -nan -gpu_avg_DC_WH, = -nan -gpu_avg_DC_WM, = -nan -gpu_avg_TC_H, = -nan -gpu_avg_TC_M, = -nan -gpu_avg_CC_H, = -nan -gpu_avg_CC_M, = -nan -gpu_avg_SHRD_ACC, = -nan -gpu_avg_REG_RD, = -nan -gpu_avg_REG_WR, = -nan -gpu_avg_NON_REG_OPs, = -nan -gpu_avg_SP_ACC, = -nan -gpu_avg_SFU_ACC, = -nan -gpu_avg_FPU_ACC, = -nan -gpu_avg_MEM_RD, = -nan -gpu_avg_MEM_WR, = -nan -gpu_avg_MEM_PRE, = -nan -gpu_avg_L2_RH, = -nan -gpu_avg_L2_RM, = -nan -gpu_avg_L2_WH, = -nan -gpu_avg_L2_WM, = -nan -gpu_avg_NOC_A, = -nan -gpu_avg_PIPE_A, = -nan -gpu_avg_IDLE_CORE_N, = -nan -gpu_avg_CONST_DYNAMICN = -nan - -Kernel Maximum Power Data: -kernel_max_power = 0 -gpu_max_IBP, = 0 -gpu_max_ICP, = 0 -gpu_max_DCP, = 0 -gpu_max_TCP, = 0 -gpu_max_CCP, = 0 -gpu_max_SHRDP, = 0 -gpu_max_RFP, = 0 -gpu_max_SPP, = 0 -gpu_max_SFUP, = 0 -gpu_max_FPUP, = 0 -gpu_max_SCHEDP, = 0 -gpu_max_L2CP, = 0 -gpu_max_MCP, = 0 -gpu_max_NOCP, = 0 -gpu_max_DRAMP, = 0 -gpu_max_PIPEP, = 0 -gpu_max_IDLE_COREP, = 0 -gpu_max_CONST_DYNAMICP = 0 -gpu_max_TOT_INST, = 0 -gpu_max_FP_INT, = 0 -gpu_max_IC_H, = 0 -gpu_max_IC_M, = 0 -gpu_max_DC_RH, = 0 -gpu_max_DC_RM, = 0 -gpu_max_DC_WH, = 0 -gpu_max_DC_WM, = 0 -gpu_max_TC_H, = 0 -gpu_max_TC_M, = 0 -gpu_max_CC_H, = 0 -gpu_max_CC_M, = 0 -gpu_max_SHRD_ACC, = 0 -gpu_max_REG_RD, = 0 -gpu_max_REG_WR, = 0 -gpu_max_NON_REG_OPs, = 0 -gpu_max_SP_ACC, = 0 -gpu_max_SFU_ACC, = 0 -gpu_max_FPU_ACC, = 0 -gpu_max_MEM_RD, = 0 -gpu_max_MEM_WR, = 0 -gpu_max_MEM_PRE, = 0 -gpu_max_L2_RH, = 0 -gpu_max_L2_RM, = 0 -gpu_max_L2_WH, = 0 -gpu_max_L2_WM, = 0 -gpu_max_NOC_A, = 0 -gpu_max_PIPE_A, = 0 -gpu_max_IDLE_CORE_N, = 0 -gpu_max_CONST_DYNAMICN = 0 - -Kernel Minimum Power Data: -kernel_min_power = 0 -gpu_min_IBP, = 0 -gpu_min_ICP, = 0 -gpu_min_DCP, = 0 -gpu_min_TCP, = 0 -gpu_min_CCP, = 0 -gpu_min_SHRDP, = 0 -gpu_min_RFP, = 0 -gpu_min_SPP, = 0 -gpu_min_SFUP, = 0 -gpu_min_FPUP, = 0 -gpu_min_SCHEDP, = 0 -gpu_min_L2CP, = 0 -gpu_min_MCP, = 0 -gpu_min_NOCP, = 0 -gpu_min_DRAMP, = 0 -gpu_min_PIPEP, = 0 -gpu_min_IDLE_COREP, = 0 -gpu_min_CONST_DYNAMICP = 0 -gpu_min_TOT_INST, = 0 -gpu_min_FP_INT, = 0 -gpu_min_IC_H, = 0 -gpu_min_IC_M, = 0 -gpu_min_DC_RH, = 0 -gpu_min_DC_RM, = 0 -gpu_min_DC_WH, = 0 -gpu_min_DC_WM, = 0 -gpu_min_TC_H, = 0 -gpu_min_TC_M, = 0 -gpu_min_CC_H, = 0 -gpu_min_CC_M, = 0 -gpu_min_SHRD_ACC, = 0 -gpu_min_REG_RD, = 0 -gpu_min_REG_WR, = 0 -gpu_min_NON_REG_OPs, = 0 -gpu_min_SP_ACC, = 0 -gpu_min_SFU_ACC, = 0 -gpu_min_FPU_ACC, = 0 -gpu_min_MEM_RD, = 0 -gpu_min_MEM_WR, = 0 -gpu_min_MEM_PRE, = 0 -gpu_min_L2_RH, = 0 -gpu_min_L2_RM, = 0 -gpu_min_L2_WH, = 0 -gpu_min_L2_WM, = 0 -gpu_min_NOC_A, = 0 -gpu_min_PIPE_A, = 0 -gpu_min_IDLE_CORE_N, = 0 -gpu_min_CONST_DYNAMICN = 0 - -Accumulative Power Statistics Over Previous Kernels: -gpu_tot_avg_power = -nan -gpu_tot_max_power = 0 -gpu_tot_min_power = 0 - - diff --git a/cuda-kernels/log b/cuda-kernels/log deleted file mode 100644 index 98df26a..0000000 --- a/cuda-kernels/log +++ /dev/null @@ -1,6328 +0,0 @@ - - - *** GPGPU-Sim Simulator Version 3.2.2 [build gpgpu-sim_git-commit-03de5ae03420ba5666d669c6f76faccf2704fa58_modified_17] *** - - -GPGPU-Sim PTX: simulation mode 0 (can change with PTX_SIM_MODE_FUNC environment variable: - 1=functional simulation only, 0=detailed performance simulator) -GPGPU-Sim: Configuration options: - --network_mode 1 # Interconnection network mode --inter_config_file config_fermi_islip.icnt # Interconnection network config file --gpgpu_ptx_use_cuobjdump 1 # Use cuobjdump to extract ptx and sass from binaries --gpgpu_experimental_lib_support 0 # Try to extract code from cuda libraries [Broken because of unknown cudaGetExportTable] --gpgpu_ptx_convert_to_ptxplus 0 # Convert SASS (native ISA) to ptxplus and run ptxplus --gpgpu_ptx_force_max_capability 70 # Force maximum compute capability --gpgpu_ptx_inst_debug_to_file 0 # Dump executed instructions' debug information to file --gpgpu_ptx_inst_debug_file inst_debug.txt # Executed instructions' debug output file --gpgpu_ptx_inst_debug_thread_uid 1 # Thread UID for executed instructions' debug output --gpgpu_simd_model 1 # 1 = post-dominator --gpgpu_shader_core_pipeline 2048:32 # shader core pipeline config, i.e., {:} --gpgpu_tex_cache:l1 16:128:24,L:R:m:N:L,F:128:4,128:2 # per-shader L1 texture cache (READ-ONLY) config {::,:::,::,:} --gpgpu_const_cache:l1 128:64:2,L:R:f:N:L,A:2:64,4 # per-shader L1 constant memory cache (READ-ONLY) config {::,:::,::,} --gpgpu_cache:il1 8:128:4,L:R:f:N:L,A:2:48,4 # shader L1 instruction cache config {::,:::,::,} --gpgpu_cache:dl1 64:128:6,L:L:m:N:H,A:128:8,8 # per-shader L1 data cache config {::,:::,::, | none} --gpgpu_cache:dl1PrefL1 none # per-shader L1 data cache config {::,:::,::, | none} --gpgpu_cache:dl1PreShared none # per-shader L1 data cache config {::,:::,::, | none} --gmem_skip_L1D 1 # global memory access skip L1D cache (implements -Xptxas -dlcm=cg, default=no skip) --gpgpu_perfect_mem 0 # enable perfect memory mode (no cache miss) --n_regfile_gating_group 4 # group of lanes that should be read/written together) --gpgpu_clock_gated_reg_file 0 # enable clock gated reg file for power calculations --gpgpu_clock_gated_lanes 0 # enable clock gated lanes for power calculations --gpgpu_shader_registers 65536 # Number of registers per shader core. Limits number of concurrent CTAs. (default 8192) --gpgpu_shader_cta 32 # Maximum number of concurrent CTAs in shader (default 8) --gpgpu_num_cta_barriers 16 # Maximum number of named barriers per CTA (default 16) --gpgpu_n_clusters 40 # number of processing clusters --gpgpu_n_cores_per_cluster 1 # number of simd cores per cluster --gpgpu_n_cluster_ejection_buffer_size 8 # number of packets in ejection buffer --gpgpu_n_ldst_response_buffer_size 2 # number of response packets in ld/st unit ejection buffer --gpgpu_shmem_size 16384 # Size of shared memory per shader core (default 16kB) --gpgpu_shmem_size 98304 # Size of shared memory per shader core (default 16kB) --gpgpu_shmem_size_PrefL1 16384 # Size of shared memory per shader core (default 16kB) --gpgpu_shmem_size_PrefShared 16384 # Size of shared memory per shader core (default 16kB) --gpgpu_shmem_num_banks 32 # Number of banks in the shared memory in each shader core (default 16) --gpgpu_shmem_limited_broadcast 0 # Limit shared memory to do one broadcast per cycle (default on) --gpgpu_shmem_warp_parts 1 # Number of portions a warp is divided into for shared memory bank conflict check --gpgpu_warpdistro_shader -1 # Specify which shader core to collect the warp size distribution from --gpgpu_warp_issue_shader 0 # Specify which shader core to collect the warp issue distribution from --gpgpu_local_mem_map 1 # Mapping from local memory space address to simulated GPU physical address space (default = enabled) --gpgpu_num_reg_banks 32 # Number of register banks (default = 8) --gpgpu_reg_bank_use_warp_id 0 # Use warp ID in mapping registers to banks (default = off) --gpgpu_operand_collector_num_units_sp 20 # number of collector units (default = 4) --gpgpu_operand_collector_num_units_sfu 4 # number of collector units (default = 4) --gpgpu_operand_collector_num_units_mem 8 # number of collector units (default = 2) --gpgpu_operand_collector_num_units_gen 0 # number of collector units (default = 0) --gpgpu_operand_collector_num_in_ports_sp 4 # number of collector unit in ports (default = 1) --gpgpu_operand_collector_num_in_ports_sfu 1 # number of collector unit in ports (default = 1) --gpgpu_operand_collector_num_in_ports_mem 1 # number of collector unit in ports (default = 1) --gpgpu_operand_collector_num_in_ports_gen 0 # number of collector unit in ports (default = 0) --gpgpu_operand_collector_num_out_ports_sp 4 # number of collector unit in ports (default = 1) --gpgpu_operand_collector_num_out_ports_sfu 1 # number of collector unit in ports (default = 1) --gpgpu_operand_collector_num_out_ports_mem 1 # number of collector unit in ports (default = 1) --gpgpu_operand_collector_num_out_ports_gen 0 # number of collector unit in ports (default = 0) --gpgpu_coalesce_arch 13 # Coalescing arch (default = 13, anything else is off for now) --gpgpu_num_sched_per_core 2 # Number of warp schedulers per core --gpgpu_max_insn_issue_per_warp 2 # Max number of instructions that can be issued per warp in one cycle by scheduler --gpgpu_simt_core_sim_order 1 # Select the simulation order of cores in a cluster (0=Fix, 1=Round-Robin) --gpgpu_pipeline_widths 4,1,1,4,1,1,6 # Pipeline widths ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB --gpgpu_num_sp_units 4 # Number of SP units (default=1) --gpgpu_num_sfu_units 1 # Number of SF units (default=1) --gpgpu_num_mem_units 1 # Number if ldst units (default=1) WARNING: not hooked up to anything --gpgpu_scheduler gto # Scheduler configuration: < lrr | gto | two_level_active > If two_level_active:::For complete list of prioritization values see shader.h enum scheduler_prioritization_typeDefault: gto --gpgpu_concurrent_kernel_sm 0 # Support concurrent kernels on a SM (default = disabled) --gpgpu_dram_scheduler 1 # 0 = fifo, 1 = FR-FCFS (defaul) --gpgpu_dram_partition_queues 8:8:8:8 # i2$:$2d:d2$:$2i --l2_ideal 0 # Use a ideal L2 cache that always hit --gpgpu_cache:dl2 64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # unified banked L2 data cache config {::,:::,::,} --gpgpu_cache:dl2_texture_only 0 # L2 cache used for texture only --gpgpu_n_mem 11 # number of memory modules (e.g. memory controllers) in gpu --gpgpu_n_sub_partition_per_mchannel 2 # number of memory subpartition in each memory module --gpgpu_n_mem_per_ctrlr 1 # number of memory chips per memory controller --gpgpu_memlatency_stat 14 # track and display latency statistics 0x2 enables MC, 0x4 enables queue logs --gpgpu_frfcfs_dram_sched_queue_size 64 # 0 = unlimited (default); # entries per chip --gpgpu_dram_return_queue_size 116 # 0 = unlimited (default); # entries per chip --gpgpu_dram_buswidth 4 # default = 4 bytes (8 bytes per cycle at DDR) --gpgpu_dram_burst_length 8 # Burst length of each DRAM request (default = 4 data bus cycle) --dram_data_command_freq_ratio 4 # Frequency ratio between DRAM data bus and command bus (default = 2 times, i.e. DDR) --gpgpu_dram_timing_opt nbk=16:CCD=2:RRD=6:RCD=12:RAS=28:RP=12:RC=40: CL=12:WL=4:CDLR=5:WR=12:nbkgrp=1:CCDL=0:RTPL=0 # DRAM timing parameters = {nbk:tCCD:tRRD:tRCD:tRAS:tRP:tRC:CL:WL:tCDLR:tWR:nbkgrp:tCCDL:tRTPL} --rop_latency 120 # ROP queue latency (default 85) --dram_latency 100 # DRAM latency (default 30) --gpgpu_mem_addr_mapping dramid@8;00000000.00000000.00000000.00000000.0000RRRR.RRRRRRRR.RBBBCCCC.BCCSSSSS # mapping memory address to dram model {dramid@;} --gpgpu_mem_addr_test 0 # run sweep test to check address mapping for aliased address --gpgpu_mem_address_mask 1 # 0 = old addressing mask, 1 = new addressing mask, 2 = new add. mask + flipped bank sel and chip sel bits --gpuwattch_xml_file gpuwattch_gtx1080Ti.xml # GPUWattch XML file --power_simulation_enabled 1 # Turn on power simulator (1=On, 0=Off) --power_per_cycle_dump 0 # Dump detailed power output each cycle --power_trace_enabled 0 # produce a file for the power trace (1=On, 0=Off) --power_trace_zlevel 6 # Compression level of the power trace output log (0=no comp, 9=highest) --steady_power_levels_enabled 0 # produce a file for the steady power levels (1=On, 0=Off) --steady_state_definition 8:4 # allowed deviation:number of samples --gpgpu_max_cycle 0 # terminates gpu simulation early (0 = no limit) --gpgpu_max_insn 0 # terminates gpu simulation early (0 = no limit) --gpgpu_max_cta 0 # terminates gpu simulation early (0 = no limit) --gpgpu_runtime_stat 500 # display runtime statistics such as dram utilization {:} --liveness_message_freq 1 # Minimum number of seconds between simulation liveness messages (0 = always print) --gpgpu_flush_l1_cache 0 # Flush L1 cache at the end of each kernel call --gpgpu_flush_l2_cache 0 # Flush L2 cache at the end of each kernel call --gpgpu_deadlock_detect 1 # Stop the simulation at deadlock (1=on (default), 0=off) --gpgpu_ptx_instruction_classification 0 # if enabled will classify ptx instruction types per kernel (Max 255 kernels now) --gpgpu_ptx_sim_mode 0 # Select between Performance (default) or Functional simulation (1) --gpgpu_clock_domains 1481.0:2962.0:1481.0:2750.0 # Clock Domain Frequencies in MhZ {:::} --gpgpu_max_concurrent_kernel 8 # maximum kernels that can run concurrently on GPU --gpgpu_cflog_interval 0 # Interval between each snapshot in control flow logger --visualizer_enabled 0 # Turn on visualizer output (1=On, 0=Off) --visualizer_outputfile NULL # Specifies the output log file for visualizer --visualizer_zlevel 6 # Compression level of the visualizer output log (0=no comp, 9=highest) --trace_enabled 0 # Turn on traces --trace_components none # comma seperated list of traces to enable. Complete list found in trace_streams.tup. Default none --trace_sampling_core 0 # The core which is printed using CORE_DPRINTF. Default 0 --trace_sampling_memory_partition -1 # The memory partition which is printed using MEMPART_DPRINTF. Default -1 (i.e. all) --enable_ptx_file_line_stats 1 # Turn on PTX source line statistic profiling. (1 = On) --ptx_line_stats_filename gpgpu_inst_stats.txt # Output file for PTX source line statistics. --gpgpu_kernel_launch_latency 0 # Kernel launch latency in cycles. Default: 0 --gpgpu_cdp_enabled 0 # Turn on CDP --save_embedded_ptx 0 # saves ptx files embedded in binary as .ptx --keep 0 # keep intermediate files created by GPGPU-Sim when interfacing with external programs --gpgpu_ptx_save_converted_ptxplus 0 # Saved converted ptxplus to a file --ptx_opcode_latency_int 4,13,4,5,145,16,4 # Opcode latencies for integers Default 1,1,19,25,145,1,4 --ptx_opcode_latency_fp 4,13,4,5,39 # Opcode latencies for single precision floating points Default 1,1,1,1,30 --ptx_opcode_latency_dp 8,19,8,8,330 # Opcode latencies for double precision floating points Default 8,8,8,8,335 --ptx_opcode_initiation_int 1,2,2,2,8,16,4 # Opcode initiation intervals for integers Default 1,1,4,4,32,1,1 --ptx_opcode_initiation_fp 1,2,1,1,4 # Opcode initiation intervals for single precision floating points Default 1,1,1,1,5 --ptx_opcode_initiation_dp 1,2,1,1,130 # Opcode initiation intervals for double precision floating points Default 8,8,8,8,130 --cdp_latency 7200,8000,100,12000,1600 # CDP API latency Default 7200,8000,100,12000,1600 -DRAM Timing Options: -nbk 16 # number of banks -CCD 2 # column to column delay -RRD 6 # minimal delay between activation of rows in different banks -RCD 12 # row to column delay -RAS 28 # time needed to activate row -RP 12 # time needed to precharge (deactivate) row -RC 40 # row cycle time -CDLR 5 # switching from write to read (changes tWTR) -WR 12 # last data-in to row precharge -CL 12 # CAS latency -WL 4 # Write latency -nbkgrp 1 # number of bank groups -CCDL 0 # column to column delay between accesses to different bank groups -RTPL 0 # read to precharge delay between accesses to different bank groups -Total number of memory sub partition = 22 -addr_dec_mask[CHIP] = 0000000000000000 high:64 low:0 -addr_dec_mask[BK] = 0000000000007080 high:15 low:7 -addr_dec_mask[ROW] = 000000000fff8000 high:28 low:15 -addr_dec_mask[COL] = 0000000000000f7f high:12 low:0 -addr_dec_mask[BURST] = 000000000000001f high:5 low:0 -sub_partition_id_mask = 0000000000000080 -GPGPU-Sim uArch: clock freqs: 1481000000.000000:2962000000.000000:1481000000.000000:2750000000.000000 -GPGPU-Sim uArch: clock periods: 0.00000000067521944632:0.00000000033760972316:0.00000000067521944632:0.00000000036363636364 -*** Initializing Memory Statistics *** -GPGPU-Sim uArch: interconnect node map (shaderID+MemID to icntID) -GPGPU-Sim uArch: Memory nodes ID start from index: 40 -GPGPU-Sim uArch: 0 1 2 3 4 5 6 -GPGPU-Sim uArch: 7 8 9 10 11 12 13 -GPGPU-Sim uArch: 14 15 16 17 18 19 20 -GPGPU-Sim uArch: 21 22 23 24 25 26 27 -GPGPU-Sim uArch: 28 29 30 31 32 33 34 -GPGPU-Sim uArch: 35 36 37 38 39 40 41 -GPGPU-Sim uArch: 42 43 44 45 46 47 48 -GPGPU-Sim uArch: 49 50 51 52 53 54 55 -GPGPU-Sim uArch: 56 57 58 59 60 61 -GPGPU-Sim uArch: interconnect node reverse map (icntID to shaderID+MemID) -GPGPU-Sim uArch: Memory nodes start from ID: 40 -GPGPU-Sim uArch: 0 1 2 3 4 5 6 -GPGPU-Sim uArch: 7 8 9 10 11 12 13 -GPGPU-Sim uArch: 14 15 16 17 18 19 20 -GPGPU-Sim uArch: 21 22 23 24 25 26 27 -GPGPU-Sim uArch: 28 29 30 31 32 33 34 -GPGPU-Sim uArch: 35 36 37 38 39 40 41 -GPGPU-Sim uArch: 42 43 44 45 46 47 48 -GPGPU-Sim uArch: 49 50 51 52 53 54 55 -GPGPU-Sim uArch: 56 57 58 59 60 61 -a9478053306cbb4803bedf0d6ea12100 /home/araihan/gpgpusim-tensorcore/cuda-kernels/tensor_core -GPGPU-Sim uArch: performance model initialization complete. -GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = 1, filename=default -self exe links to: /home/araihan/gpgpusim-tensorcore/cuda-kernels/tensor_core -Running md5sum using "md5sum /home/araihan/gpgpusim-tensorcore/cuda-kernels/tensor_core " -Parsing file _cuobjdump_complete_output_c7ZC8M -######### cuobjdump parser ######## -## Adding new section PTX -Adding ptx filename: _cuobjdump_1.ptx -Adding arch: sm_70 -Adding identifier: default -Done parsing!!! -GPGPU-Sim PTX: __cudaRegisterFunction _Z17convertFp32ToFp16P6__halfPfi : hostFun 0x0x401dd7, fat_cubin_handle = 1 -WARNING: No guarantee that PTX will be parsed for SM version 70 - _1.ptx:13 => (ptx_parser.cc:175) start_function - _1.ptx:13 => (ptx_parser.cc:144) init_directive_state - _1.ptx:13 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:13 => (ptx_parser.cc:144) init_directive_state - _1.ptx:13 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:13 => (ptx_parser.cc:631) add_scalar_type_spec "B32_TYPE" - _1.ptx:13 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_local scalar_type_spec=B32_TYPE - _1.ptx:14 => (ptx_parser.cc:189) add_function_name vprintf (extern) - _1.ptx:14 => (ptx_parser.cc:381) add_identifier "func_retval0" (0) -GPGPU-Sim PTX: allocating stack frame region for .param "func_retval0" from 0x0 to 0x4 - _1.ptx:14 => (ptx_parser.cc:144) init_directive_state - _1.ptx:15 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:15 => (ptx_parser.cc:631) add_scalar_type_spec "B64_TYPE" - _1.ptx:15 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_local scalar_type_spec=B64_TYPE - _1.ptx:15 => (ptx_parser.cc:381) add_identifier "vprintf_param_0" (1) -GPGPU-Sim PTX: allocating stack frame region for .param "vprintf_param_0" from 0x4 to 0xc - _1.ptx:15 => (ptx_parser.cc:577) add_function_arg "vprintf_param_0" - _1.ptx:15 => (ptx_parser.cc:219) add_directive - _1.ptx:15 => (ptx_parser.cc:144) init_directive_state - _1.ptx:16 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:16 => (ptx_parser.cc:631) add_scalar_type_spec "B64_TYPE" - _1.ptx:16 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_local scalar_type_spec=B64_TYPE - _1.ptx:17 => (ptx_parser.cc:381) add_identifier "vprintf_param_1" (2) -GPGPU-Sim PTX: allocating stack frame region for .param "vprintf_param_1" from 0xc to 0x14 - _1.ptx:17 => (ptx_parser.cc:577) add_function_arg "vprintf_param_1" - _1.ptx:17 => (ptx_parser.cc:219) add_directive - _1.ptx:17 => (ptx_parser.cc:144) init_directive_state - _1.ptx:19 => (ptx_parser.cc:605) add_space_spec "global_space" - _1.ptx:19 => (ptx_parser.cc:590) add_alignment_spec - _1.ptx:19 => (ptx_parser.cc:631) add_scalar_type_spec "B8_TYPE" - _1.ptx:19 => (ptx_parser.cc:331) set_variable_type space_spec=global_space scalar_type_spec=B8_TYPE - _1.ptx:19 => (ptx_parser.cc:381) add_identifier "$str" (3) -GPGPU-Sim PTX: allocating global region for "$str" from 0x100 to 0x109 (global memory space) - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:883) add_literal_int - _1.ptx:19 => (ptx_parser.cc:319) add_variables - _1.ptx:19 => (ptx_parser.cc:144) init_directive_state - _1.ptx:21 => (ptx_parser.cc:175) start_function - _1.ptx:21 => (ptx_parser.cc:144) init_directive_state - _1.ptx:21 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:21 => (ptx_parser.cc:144) init_directive_state - _1.ptx:21 => (ptx_parser.cc:189) add_function_name _Z12wmma_exampleP6__halfS0_Pfiiiff (entrypoint) - _1.ptx:22 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:22 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:22 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=U64_TYPE - _1.ptx:22 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0" (4) - _1.ptx:22 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0" - _1.ptx:22 => (ptx_parser.cc:219) add_directive - _1.ptx:22 => (ptx_parser.cc:144) init_directive_state - _1.ptx:23 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:23 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:23 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=U64_TYPE - _1.ptx:23 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1" (5) - _1.ptx:23 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1" - _1.ptx:23 => (ptx_parser.cc:219) add_directive - _1.ptx:23 => (ptx_parser.cc:144) init_directive_state - _1.ptx:24 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:24 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:24 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=U64_TYPE - _1.ptx:24 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2" (6) - _1.ptx:24 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2" - _1.ptx:24 => (ptx_parser.cc:219) add_directive - _1.ptx:24 => (ptx_parser.cc:144) init_directive_state - _1.ptx:25 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:25 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:25 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=U32_TYPE - _1.ptx:25 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3" (7) - _1.ptx:25 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3" - _1.ptx:25 => (ptx_parser.cc:219) add_directive - _1.ptx:25 => (ptx_parser.cc:144) init_directive_state - _1.ptx:26 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:26 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:26 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=U32_TYPE - _1.ptx:26 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4" (8) - _1.ptx:26 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4" - _1.ptx:26 => (ptx_parser.cc:219) add_directive - _1.ptx:26 => (ptx_parser.cc:144) init_directive_state - _1.ptx:27 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:27 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:27 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=U32_TYPE - _1.ptx:27 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5" (9) - _1.ptx:27 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5" - _1.ptx:27 => (ptx_parser.cc:219) add_directive - _1.ptx:27 => (ptx_parser.cc:144) init_directive_state - _1.ptx:28 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:28 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:28 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=F32_TYPE - _1.ptx:28 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_6" (10) - _1.ptx:28 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_6" - _1.ptx:28 => (ptx_parser.cc:219) add_directive - _1.ptx:28 => (ptx_parser.cc:144) init_directive_state - _1.ptx:29 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:29 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:29 => (ptx_parser.cc:331) set_variable_type space_spec=param_space_kernel scalar_type_spec=F32_TYPE - _1.ptx:30 => (ptx_parser.cc:381) add_identifier "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_7" (11) - _1.ptx:30 => (ptx_parser.cc:577) add_function_arg "_Z12wmma_exampleP6__halfS0_Pfiiiff_param_7" - _1.ptx:30 => (ptx_parser.cc:219) add_directive - _1.ptx:30 => (ptx_parser.cc:144) init_directive_state - _1.ptx:32 => (ptx_parser.cc:605) add_space_spec "local_space" - _1.ptx:32 => (ptx_parser.cc:590) add_alignment_spec - _1.ptx:32 => (ptx_parser.cc:631) add_scalar_type_spec "B8_TYPE" - _1.ptx:32 => (ptx_parser.cc:331) set_variable_type space_spec=local_space scalar_type_spec=B8_TYPE - _1.ptx:32 => (ptx_parser.cc:381) add_identifier "__local_depot0" (12) -GPGPU-Sim PTX: allocating stack frame region for .local "__local_depot0" from 0x0 to 0x8 - _1.ptx:32 => (ptx_parser.cc:319) add_variables - _1.ptx:32 => (ptx_parser.cc:144) init_directive_state - _1.ptx:32 => (ptx_parser.cc:219) add_directive - _1.ptx:32 => (ptx_parser.cc:144) init_directive_state - _1.ptx:33 => (ptx_parser.cc:605) add_space_spec "reg_space" - _1.ptx:33 => (ptx_parser.cc:631) add_scalar_type_spec "B64_TYPE" - _1.ptx:33 => (ptx_parser.cc:331) set_variable_type space_spec=reg_space scalar_type_spec=B64_TYPE - _1.ptx:33 => (ptx_parser.cc:381) add_identifier "%SP" (13) - _1.ptx:33 => (ptx_parser.cc:319) add_variables - _1.ptx:33 => (ptx_parser.cc:144) init_directive_state - _1.ptx:33 => (ptx_parser.cc:219) add_directive - _1.ptx:33 => (ptx_parser.cc:144) init_directive_state - _1.ptx:34 => (ptx_parser.cc:605) add_space_spec "reg_space" - _1.ptx:34 => (ptx_parser.cc:631) add_scalar_type_spec "B64_TYPE" - _1.ptx:34 => (ptx_parser.cc:331) set_variable_type space_spec=reg_space scalar_type_spec=B64_TYPE - _1.ptx:34 => (ptx_parser.cc:381) add_identifier "%SPL" (14) - _1.ptx:34 => (ptx_parser.cc:319) add_variables - _1.ptx:34 => (ptx_parser.cc:144) init_directive_state - _1.ptx:34 => (ptx_parser.cc:219) add_directive - _1.ptx:34 => (ptx_parser.cc:144) init_directive_state - _1.ptx:35 => (ptx_parser.cc:605) add_space_spec "reg_space" - _1.ptx:35 => (ptx_parser.cc:631) add_scalar_type_spec "PRED_TYPE" - _1.ptx:35 => (ptx_parser.cc:331) set_variable_type space_spec=reg_space scalar_type_spec=PRED_TYPE - _1.ptx:35 => (ptx_parser.cc:381) add_identifier "%p0" (15) - _1.ptx:35 => (ptx_parser.cc:381) add_identifier "%p1" (16) - _1.ptx:35 => (ptx_parser.cc:381) add_identifier "%p2" (17) - _1.ptx:35 => (ptx_parser.cc:381) add_identifier "%p3" (18) - _1.ptx:35 => (ptx_parser.cc:381) add_identifier "%p4" (19) - _1.ptx:35 => (ptx_parser.cc:381) add_identifier "%p5" (20) - _1.ptx:35 => (ptx_parser.cc:319) add_variables - _1.ptx:35 => (ptx_parser.cc:144) init_directive_state - _1.ptx:35 => (ptx_parser.cc:219) add_directive - _1.ptx:35 => (ptx_parser.cc:144) init_directive_state - _1.ptx:36 => (ptx_parser.cc:605) add_space_spec "reg_space" - _1.ptx:36 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:36 => (ptx_parser.cc:331) set_variable_type space_spec=reg_space scalar_type_spec=F32_TYPE - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f0" (21) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f1" (22) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f2" (23) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f3" (24) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f4" (25) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f5" (26) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f6" (27) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f7" (28) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f8" (29) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f9" (30) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f10" (31) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f11" (32) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f12" (33) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f13" (34) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f14" (35) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f15" (36) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f16" (37) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f17" (38) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f18" (39) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f19" (40) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f20" (41) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f21" (42) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f22" (43) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f23" (44) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f24" (45) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f25" (46) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f26" (47) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f27" (48) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f28" (49) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f29" (50) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f30" (51) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f31" (52) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f32" (53) - _1.ptx:36 => (ptx_parser.cc:381) add_identifier "%f33" (54) - _1.ptx:36 => (ptx_parser.cc:319) add_variables - _1.ptx:36 => (ptx_parser.cc:144) init_directive_state - _1.ptx:36 => (ptx_parser.cc:219) add_directive - _1.ptx:36 => (ptx_parser.cc:144) init_directive_state - _1.ptx:37 => (ptx_parser.cc:605) add_space_spec "reg_space" - _1.ptx:37 => (ptx_parser.cc:631) add_scalar_type_spec "B32_TYPE" - _1.ptx:37 => (ptx_parser.cc:331) set_variable_type space_spec=reg_space scalar_type_spec=B32_TYPE - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r0" (55) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r1" (56) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r2" (57) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r3" (58) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r4" (59) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r5" (60) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r6" (61) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r7" (62) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r8" (63) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r9" (64) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r10" (65) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r11" (66) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r12" (67) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r13" (68) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r14" (69) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r15" (70) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r16" (71) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r17" (72) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r18" (73) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r19" (74) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r20" (75) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r21" (76) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r22" (77) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r23" (78) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r24" (79) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r25" (80) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r26" (81) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r27" (82) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r28" (83) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r29" (84) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r30" (85) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r31" (86) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r32" (87) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r33" (88) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r34" (89) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r35" (90) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r36" (91) - _1.ptx:37 => (ptx_parser.cc:381) add_identifier "%r37" (92) - _1.ptx:37 => (ptx_parser.cc:319) add_variables - _1.ptx:37 => (ptx_parser.cc:144) init_directive_state - _1.ptx:37 => (ptx_parser.cc:219) add_directive - _1.ptx:37 => (ptx_parser.cc:144) init_directive_state - _1.ptx:38 => (ptx_parser.cc:605) add_space_spec "reg_space" - _1.ptx:38 => (ptx_parser.cc:631) add_scalar_type_spec "B64_TYPE" - _1.ptx:38 => (ptx_parser.cc:331) set_variable_type space_spec=reg_space scalar_type_spec=B64_TYPE - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd0" (93) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd1" (94) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd2" (95) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd3" (96) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd4" (97) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd5" (98) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd6" (99) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd7" (100) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd8" (101) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd9" (102) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd10" (103) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd11" (104) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd12" (105) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd13" (106) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd14" (107) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd15" (108) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd16" (109) - _1.ptx:38 => (ptx_parser.cc:381) add_identifier "%rd17" (110) - _1.ptx:38 => (ptx_parser.cc:319) add_variables - _1.ptx:38 => (ptx_parser.cc:144) init_directive_state - _1.ptx:38 => (ptx_parser.cc:219) add_directive - _1.ptx:38 => (ptx_parser.cc:144) init_directive_state - _1.ptx:41 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:41 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:41 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:41 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:41 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:41 => (ptx_parser.cc:144) init_directive_state - _1.ptx:42 => (ptx_parser.cc:605) add_space_spec "local_space" - _1.ptx:42 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:42 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:42 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:42 => (ptx_parser.cc:295) add_instruction: cvta - _1.ptx:42 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:42 => (ptx_parser.cc:144) init_directive_state - _1.ptx:43 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:43 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:43 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:43 => (ptx_parser.cc:929) add_address_operand - _1.ptx:43 => (ptx_parser.cc:766) add_memory_operand - _1.ptx:43 => (ptx_parser.cc:295) add_instruction: ld - _1.ptx:43 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:43 => (ptx_parser.cc:144) init_directive_state - _1.ptx:44 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:44 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:44 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:44 => (ptx_parser.cc:929) add_address_operand - _1.ptx:44 => (ptx_parser.cc:766) add_memory_operand - _1.ptx:44 => (ptx_parser.cc:295) add_instruction: ld - _1.ptx:44 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:44 => (ptx_parser.cc:144) init_directive_state - _1.ptx:45 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:45 => (ptx_parser.cc:631) add_scalar_type_spec "U64_TYPE" - _1.ptx:45 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:45 => (ptx_parser.cc:929) add_address_operand - _1.ptx:45 => (ptx_parser.cc:766) add_memory_operand - _1.ptx:45 => (ptx_parser.cc:295) add_instruction: ld - _1.ptx:45 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:45 => (ptx_parser.cc:144) init_directive_state - _1.ptx:46 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:46 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:46 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:46 => (ptx_parser.cc:929) add_address_operand - _1.ptx:46 => (ptx_parser.cc:766) add_memory_operand - _1.ptx:46 => (ptx_parser.cc:295) add_instruction: ld - _1.ptx:46 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:46 => (ptx_parser.cc:144) init_directive_state - _1.ptx:47 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:47 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:47 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:47 => (ptx_parser.cc:929) add_address_operand - _1.ptx:47 => (ptx_parser.cc:766) add_memory_operand - _1.ptx:47 => (ptx_parser.cc:295) add_instruction: ld - _1.ptx:47 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:47 => (ptx_parser.cc:144) init_directive_state - _1.ptx:48 => (ptx_parser.cc:605) add_space_spec "param_space_unclassified" - _1.ptx:48 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:48 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:48 => (ptx_parser.cc:929) add_address_operand - _1.ptx:48 => (ptx_parser.cc:766) add_memory_operand - _1.ptx:48 => (ptx_parser.cc:295) add_instruction: ld - _1.ptx:48 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:48 => (ptx_parser.cc:144) init_directive_state - _1.ptx:50 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:50 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:50 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:50 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:50 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:50 => (ptx_parser.cc:144) init_directive_state - _1.ptx:52 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:52 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:52 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:52 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:52 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:52 => (ptx_parser.cc:144) init_directive_state - _1.ptx:53 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:53 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:53 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:53 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:53 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:53 => (ptx_parser.cc:144) init_directive_state - _1.ptx:54 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:54 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:54 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:54 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:54 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:54 => (ptx_parser.cc:144) init_directive_state - _1.ptx:55 => (ptx_parser.cc:672) add_option - _1.ptx:55 => (ptx_parser.cc:631) add_scalar_type_spec "S32_TYPE" - _1.ptx:55 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:55 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:55 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:55 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:55 => (ptx_parser.cc:295) add_instruction: mad - _1.ptx:55 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:55 => (ptx_parser.cc:144) init_directive_state - _1.ptx:56 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:56 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:56 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:56 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:56 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:56 => (ptx_parser.cc:144) init_directive_state - _1.ptx:57 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:57 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:57 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:57 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:57 => (ptx_parser.cc:295) add_instruction: div - _1.ptx:57 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:57 => (ptx_parser.cc:144) init_directive_state - _1.ptx:58 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:58 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:58 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:58 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:58 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:58 => (ptx_parser.cc:144) init_directive_state - _1.ptx:59 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:59 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:59 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:59 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:59 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:59 => (ptx_parser.cc:144) init_directive_state - _1.ptx:60 => (ptx_parser.cc:631) add_scalar_type_spec "U32_TYPE" - _1.ptx:60 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:60 => (ptx_parser.cc:760) add_builtin_operand - _1.ptx:60 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:60 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:60 => (ptx_parser.cc:144) init_directive_state - _1.ptx:61 => (ptx_parser.cc:672) add_option - _1.ptx:61 => (ptx_parser.cc:631) add_scalar_type_spec "S32_TYPE" - _1.ptx:61 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:61 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:61 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:61 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:61 => (ptx_parser.cc:295) add_instruction: mad - _1.ptx:61 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:61 => (ptx_parser.cc:144) init_directive_state - _1.ptx:62 => (ptx_parser.cc:631) add_scalar_type_spec "B32_TYPE" - _1.ptx:62 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:62 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:62 => (ptx_parser.cc:883) add_literal_int - _1.ptx:62 => (ptx_parser.cc:295) add_instruction: shl - _1.ptx:62 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:62 => (ptx_parser.cc:144) init_directive_state - _1.ptx:63 => (ptx_parser.cc:631) add_scalar_type_spec "B32_TYPE" - _1.ptx:63 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:63 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:63 => (ptx_parser.cc:883) add_literal_int - _1.ptx:63 => (ptx_parser.cc:295) add_instruction: shl - _1.ptx:63 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:63 => (ptx_parser.cc:144) init_directive_state - _1.ptx:64 => (ptx_parser.cc:672) add_option - _1.ptx:64 => (ptx_parser.cc:631) add_scalar_type_spec "S32_TYPE" - _1.ptx:64 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:64 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:64 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:64 => (ptx_parser.cc:295) add_instruction: setp - _1.ptx:64 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:64 => (ptx_parser.cc:144) init_directive_state - _1.ptx:65 => (ptx_parser.cc:672) add_option - _1.ptx:65 => (ptx_parser.cc:631) add_scalar_type_spec "S32_TYPE" - _1.ptx:65 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:65 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:65 => (ptx_parser.cc:883) add_literal_int - _1.ptx:65 => (ptx_parser.cc:295) add_instruction: setp - _1.ptx:65 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:65 => (ptx_parser.cc:144) init_directive_state - _1.ptx:66 => (ptx_parser.cc:631) add_scalar_type_spec "PRED_TYPE" - _1.ptx:66 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:66 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:66 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:66 => (ptx_parser.cc:295) add_instruction: and - _1.ptx:66 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:66 => (ptx_parser.cc:144) init_directive_state - _1.ptx:67 => (ptx_parser.cc:672) add_option - _1.ptx:67 => (ptx_parser.cc:631) add_scalar_type_spec "S32_TYPE" - _1.ptx:67 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:67 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:67 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:67 => (ptx_parser.cc:295) add_instruction: setp - _1.ptx:67 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:67 => (ptx_parser.cc:144) init_directive_state - _1.ptx:68 => (ptx_parser.cc:631) add_scalar_type_spec "PRED_TYPE" - _1.ptx:68 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:68 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:68 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:68 => (ptx_parser.cc:295) add_instruction: and - _1.ptx:68 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:68 => (ptx_parser.cc:144) init_directive_state - _1.ptx:69 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:69 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:69 => (ptx_parser.cc:889) add_literal_float - _1.ptx:69 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:69 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:69 => (ptx_parser.cc:144) init_directive_state - _1.ptx:70 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:70 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:70 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:70 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:70 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:70 => (ptx_parser.cc:144) init_directive_state - _1.ptx:71 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:71 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:71 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:71 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:71 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:71 => (ptx_parser.cc:144) init_directive_state - _1.ptx:72 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:72 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:72 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:72 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:72 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:72 => (ptx_parser.cc:144) init_directive_state - _1.ptx:73 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:73 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:73 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:73 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:73 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:73 => (ptx_parser.cc:144) init_directive_state - _1.ptx:74 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:74 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:74 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:74 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:74 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:74 => (ptx_parser.cc:144) init_directive_state - _1.ptx:75 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:75 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:75 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:75 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:75 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:75 => (ptx_parser.cc:144) init_directive_state - _1.ptx:76 => (ptx_parser.cc:631) add_scalar_type_spec "F32_TYPE" - _1.ptx:76 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:76 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:76 => (ptx_parser.cc:295) add_instruction: mov - _1.ptx:76 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:76 => (ptx_parser.cc:144) init_directive_state - _1.ptx:77 => (ptx_parser.cc:659) add_pred - _1.ptx:77 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:77 => (ptx_parser.cc:295) add_instruction: bra - _1.ptx:77 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:77 => (ptx_parser.cc:144) init_directive_state - _1.ptx:78 => (ptx_parser.cc:672) add_option - _1.ptx:78 => (ptx_parser.cc:901) add_scalar_operand - _1.ptx:78 => (ptx_parser.cc:295) add_instruction: bra - _1.ptx:78 => (ptx_parser.cc:159) init_instruction_state - _1.ptx:78 => (ptx_parser.cc:144) init_directive_state - _1.ptx:80 => (ptx_parser.cc:643) add_label - _1.ptx:80 => (ptx_parser.cc:295) add_instruction: