summaryrefslogtreecommitdiff
path: root/cuda-kernels/_cuobjdump_complete_output_EIGzTK
diff options
context:
space:
mode:
Diffstat (limited to 'cuda-kernels/_cuobjdump_complete_output_EIGzTK')
-rw-r--r--cuda-kernels/_cuobjdump_complete_output_EIGzTK1055
1 files changed, 1055 insertions, 0 deletions
diff --git a/cuda-kernels/_cuobjdump_complete_output_EIGzTK b/cuda-kernels/_cuobjdump_complete_output_EIGzTK
new file mode 100644
index 0000000..36999c0
--- /dev/null
+++ b/cuda-kernels/_cuobjdump_complete_output_EIGzTK
@@ -0,0 +1,1055 @@
+
+Fatbin elf code:
+================
+arch = sm_70
+code version = [1,7]
+producer = <unknown>
+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;
+}
+
+
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:35:00 +0000