diff options
26 files changed, 4918 insertions, 9 deletions
diff --git a/cuda-kernels/Makefile b/cuda-kernels/Makefile new file mode 100755 index 0000000..51a7760 --- /dev/null +++ b/cuda-kernels/Makefile @@ -0,0 +1,7 @@ +all: tensor_core.cu + nvcc -arch=sm_70 -lcudart -g -o tensor_core tensor_core.cu + +.PHONY: +clean: + rm tensorcore +# nvcc -arch=sm_70 --gpu-architecture=compute_50 --gpu-code=compute_50 -lcudart -g -o tensor_core tensor_core.cu diff --git a/cuda-kernels/_cuobjdump_1.elf b/cuda-kernels/_cuobjdump_1.elf new file mode 100644 index 0000000..672b0f0 --- /dev/null +++ b/cuda-kernels/_cuobjdump_1.elf @@ -0,0 +1,15 @@ +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 new file mode 100644 index 0000000..3453f4a --- /dev/null +++ b/cuda-kernels/_cuobjdump_1.ptx @@ -0,0 +1,170 @@ + + + + + + + +.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 new file mode 100644 index 0000000..2aac29a --- /dev/null +++ b/cuda-kernels/_cuobjdump_1.sass @@ -0,0 +1,2 @@ + code for sm_70 + diff --git a/cuda-kernels/_cuobjdump_2.elf b/cuda-kernels/_cuobjdump_2.elf new file mode 100644 index 0000000..c03b06d --- /dev/null +++ b/cuda-kernels/_cuobjdump_2.elf @@ -0,0 +1,494 @@ +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 new file mode 100644 index 0000000..1b50ed2 --- /dev/null +++ b/cuda-kernels/_cuobjdump_2.sass @@ -0,0 +1,348 @@ + 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 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; +} + + diff --git a/cuda-kernels/_cuobjdump_complete_output_rndQyq b/cuda-kernels/_cuobjdump_complete_output_rndQyq new file mode 100644 index 0000000..36999c0 --- /dev/null +++ b/cuda-kernels/_cuobjdump_complete_output_rndQyq @@ -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; +} + + diff --git a/cuda-kernels/config_fermi_islip.icnt b/cuda-kernels/config_fermi_islip.icnt new file mode 100755 index 0000000..a788090 --- /dev/null +++ b/cuda-kernels/config_fermi_islip.icnt @@ -0,0 +1,70 @@ +//21*1 fly with 32 flits per packet under gpgpusim injection mode +use_map = 0; +flit_size = 32; + +// currently we do not use this, see subnets below +network_count = 2; + +// Topology +topology = fly; +k = 62; +n = 1; + +// Routing + +routing_function = dest_tag; + +// Flow control + +num_vcs = 1; +vc_buf_size = 8; + +wait_for_tail_credit = 0; + +// Router architecture + +vc_allocator = islip; //separable_input_first; +sw_allocator = islip; //separable_input_first; +alloc_iters = 1; + +credit_delay = 0; +routing_delay = 0; +vc_alloc_delay = 1; +sw_alloc_delay = 1; + +input_speedup = 2; +output_speedup = 1; +internal_speedup = 1.0; + +// Traffic, GPGPU-Sim does not use this + +traffic = uniform; +packet_size ={{1,2,3,4},{10,20}}; +packet_size_rate={{1,1,1,1},{2,1}}; + +// Simulation - Don't change + +sim_type = gpgpusim; +//sim_type = latency; +injection_rate = 0.1; + +subnets = 2; + +// Always use read and write no matter following line +//use_read_write = 1; + + +read_request_subnet = 0; +read_reply_subnet = 1; +write_request_subnet = 0; +write_reply_subnet = 1; + +read_request_begin_vc = 0; +read_request_end_vc = 0; +write_request_begin_vc = 0; +write_request_end_vc = 0; +read_reply_begin_vc = 0; +read_reply_end_vc = 0; +write_reply_begin_vc = 0; +write_reply_end_vc = 0; + diff --git a/cuda-kernels/gpgpu_inst_stats.txt b/cuda-kernels/gpgpu_inst_stats.txt new file mode 100755 index 0000000..acb1839 --- /dev/null +++ b/cuda-kernels/gpgpu_inst_stats.txt @@ -0,0 +1 @@ +kernel line : count latency dram_traffic smem_bk_conflicts smem_warp gmem_access_generated gmem_warp exposed_latency warp_divergence diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config new file mode 100755 index 0000000..306d7f9 --- /dev/null +++ b/cuda-kernels/gpgpusim.config @@ -0,0 +1,149 @@ +# This config models the Pascal GP102 (GeForceGTX 1080Ti) + +# functional simulator specification +-gpgpu_ptx_instruction_classification 0 +-gpgpu_ptx_sim_mode 0 +-gpgpu_ptx_force_max_capability 70 + +# SASS execution (only supported with CUDA >= 4.0) +-gpgpu_ptx_convert_to_ptxplus 0 +-gpgpu_ptx_save_converted_ptxplus 0 + +# high level architecture configuration +-gpgpu_n_clusters 40 +-gpgpu_n_cores_per_cluster 1 +-gpgpu_n_mem 11 +-gpgpu_n_sub_partition_per_mchannel 2 + +# Pascal clock domains +#-gpgpu_clock_domains <Core Clock>:<Interconnect Clock>:<L2 Clock>:<DRAM Clock> +# Pascal NVIDIA TITAN X clock domains are adopted from +# https://en.wikipedia.org/wiki/GeForce_10_series +-gpgpu_clock_domains 1481.0:2962.0:1481.0:2750.0 + +# shader core pipeline config +-gpgpu_shader_registers 65536 + +# This implies a maximum of 64 warps/SM +-gpgpu_shader_core_pipeline 2048:32 +-gpgpu_shader_cta 32 +-gpgpu_simd_model 1 + +# Pipeline widths and number of FUs +# ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB +## Pascal GP102 has 4 SP SIMD units and 1 SFU unit +## we need to scale the number of pipeline registers to be equal to the number of SP units +-gpgpu_pipeline_widths 4,1,1,4,1,1,6 +-gpgpu_num_sp_units 4 +-gpgpu_num_sfu_units 1 + +# Instruction latencies and initiation intervals +# "ADD,MAX,MUL,MAD,DIV" +# SFU is 32-width in pascal, then dp units initiation is 1 cycle +-ptx_opcode_latency_int 4,13,4,5,145,16,4 +-ptx_opcode_initiation_int 1,2,2,2,8,16,4 +-ptx_opcode_latency_fp 4,13,4,5,39 +-ptx_opcode_initiation_fp 1,2,1,1,4 +-ptx_opcode_latency_dp 8,19,8,8,330 +-ptx_opcode_initiation_dp 1,2,1,1,130 + +# <nsets>:<bsize>:<assoc>,<rep>:<wr>:<alloc>:<wr_alloc>:<set_index_fn>,<mshr>:<N>:<merge>,<mq>:**<fifo_entry> +# ** Optional parameter - Required when mshr_type==Texture Fifo +# Note: Hashing set index function (H) only applies to a set size of 32 or 64. +# Pascal GP102 has 96KB Shared memory +# Pascal GP102 has 64KB L1 cache +# The default is to disable the L1 cache, unless cache modifieres is used +-gpgpu_cache:dl1 64:128:6,L:L:m:N:H,A:128:8,8 +-gpgpu_shmem_size 98304 +-gmem_skip_L1D 1 + +# 64 sets, each 128 bytes 16-way for each memory sub partition (128 KB per memory sub partition). This gives 3MB L2 cache +-gpgpu_cache:dl2 64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # used to be 128:4 +-gpgpu_cache:dl2_texture_only 0 + +# 4 KB Inst. +-gpgpu_cache:il1 8:128:4,L:R:f:N:L,A:2:48,4 +# 48 KB Tex +-gpgpu_tex_cache:l1 16:128:24,L:R:m:N:L,F:128:4,128:2 +# 12 KB Const +-gpgpu_const_cache:l1 128:64:2,L:R:f:N:L,A:2:64,4 + +# enable operand collector +## larger operand collectors and reg_banks are needed for the 4 warp schedulers and 4 SIMD units +-gpgpu_operand_collector_num_units_sp 20 +-gpgpu_operand_collector_num_units_sfu 4 +-gpgpu_operand_collector_num_units_mem 8 +-gpgpu_operand_collector_num_in_ports_sp 4 +-gpgpu_operand_collector_num_out_ports_sp 4 +-gpgpu_operand_collector_num_in_ports_sfu 1 +-gpgpu_operand_collector_num_out_ports_sfu 1 +-gpgpu_operand_collector_num_in_ports_mem 1 +-gpgpu_operand_collector_num_out_ports_mem 1 +# gpgpu_num_reg_banks should be increased to 32, but it gives an error! +-gpgpu_num_reg_banks 32 + +# shared memory bankconflict detection +-gpgpu_shmem_num_banks 32 +-gpgpu_shmem_limited_broadcast 0 +-gpgpu_shmem_warp_parts 1 + +## In Pascal, a warp scheduler can issue 2 insts per cycle +-gpgpu_max_insn_issue_per_warp 2 + +# interconnection +-network_mode 1 +-inter_config_file config_fermi_islip.icnt + +# memory partition latency config +-rop_latency 120 +-dram_latency 100 + +# dram model config +-gpgpu_dram_scheduler 1 +# The DRAM return queue and the scheduler queue together should provide buffer +# to sustain the memory level parallelism to tolerate DRAM latency +# To allow 100% DRAM utility, there should at least be enough buffer to sustain +# the minimum DRAM latency (100 core cycles). I.e. +# Total buffer space required = 100 x 924MHz / 700MHz = 132 +-gpgpu_frfcfs_dram_sched_queue_size 64 +-gpgpu_dram_return_queue_size 116 + +# for NVIDIA GeForceGTX 1080Ti, bus width is 352bits (11 DRAM chips x 32 bits) +# 11 memory paritions, 4 bytes (1 DRAM chip) per memory partition +# the atom size of GDDR5X (the smallest read request) is 32 bytes +-gpgpu_n_mem_per_ctrlr 1 +-gpgpu_dram_buswidth 4 +-gpgpu_dram_burst_length 8 +-dram_data_command_freq_ratio 4 # GDDR5X is QDR +-gpgpu_mem_address_mask 1 +-gpgpu_mem_addr_mapping dramid@8;00000000.00000000.00000000.00000000.0000RRRR.RRRRRRRR.RBBBCCCC.BCCSSSSS + +# Use the same GDDR5 timing from hynix H5GQ1H24AFR +# disable bank groups for now, set nbkgrp to 1 and tCCDL and tRTPL to 0 +-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" + +# Pascal has four schedulers per core +-gpgpu_num_sched_per_core 2 +# Two Level Scheduler with active and pending pools +#-gpgpu_scheduler two_level_active:6:0:1 +# Loose round robbin scheduler +#-gpgpu_scheduler lrr +# Greedy then oldest scheduler +-gpgpu_scheduler gto + +# stat collection +-gpgpu_memlatency_stat 14 +-gpgpu_runtime_stat 500 +-enable_ptx_file_line_stats 1 +-visualizer_enabled 0 + +# power model configs +-power_simulation_enabled 1 +-gpuwattch_xml_file gpuwattch_gtx1080Ti.xml + +# tracing functionality +#-trace_enabled 1 +#-trace_components WARP_SCHEDULER,SCOREBOARD +#-trace_sampling_core 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 new file mode 100644 index 0000000..f754f0c --- /dev/null +++ b/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-34-2018.log @@ -0,0 +1,324 @@ +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 new file mode 100644 index 0000000..f754f0c --- /dev/null +++ b/cuda-kernels/gpgpusim_power_report__Sun-May-27-14-17-47-2018.log @@ -0,0 +1,324 @@ +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/gpuwattch_gtx1080Ti.xml b/cuda-kernels/gpuwattch_gtx1080Ti.xml new file mode 100755 index 0000000..02619ff --- /dev/null +++ b/cuda-kernels/gpuwattch_gtx1080Ti.xml @@ -0,0 +1,538 @@ +<?xml version="1.0" ?> +<component id="root" name="root"> + <component id="system" name="system"> + <!--McPAT will skip the components if number is set to 0 --> + <param name="GPU_Architecture" value="1"/><!-- 0-G80; 1-Fermi; others not supported --> + <param name="number_of_cores" value="28"/> + <param name="architecture" value="1"/> <!-- fermi:1 quadro:2 other: undefined--> + <param name="number_of_L1Directories" value="0"/> + <param name="number_of_L2Directories" value="0"/> + <param name="number_of_L2s" value="1"/> <!-- This number means how many L2 clusters in each cluster there can be multiple banks/ports --> + <param name="number_of_L3s" value="0"/> <!-- This number means how many L3 clusters --> + <param name="number_of_NoCs" value="1"/> + <param name="homogeneous_cores" value="1"/><!--1 means homo --> + <param name="homogeneous_L2s" value="1"/> + <param name="homogeneous_L1Directorys" value="1"/> + <param name="homogeneous_L2Directorys" value="1"/> + <param name="homogeneous_L3s" value="1"/> + <param name="homogeneous_ccs" value="1"/><!--cache coherece hardware --> + <param name="homogeneous_NoCs" value="1"/> + <param name="core_tech_node" value="23"/><!-- nm --> + <param name="target_core_clockrate" value="1481"/><!--MHz --> + <param name="temperature" value="380"/> <!-- Kelvin --> + <param name="number_cache_levels" value="2"/> + <param name="interconnect_projection_type" value="0"/><!--0: agressive wire technology; 1: conservative wire technology --> + <param name="device_type" value="0"/><!--0: HP(High Performance Type); 1: LSTP(Low standby power) 2: LOP (Low Operating Power) --> + <param name="longer_channel_device" value="1"/><!-- 0 no use; 1 use when possible --> + <param name="machine_bits" value="32"/> + <param name="virtual_address_width" value="32"/> + <param name="physical_address_width" value="32"/> + <param name="virtual_memory_page_size" value="4096"/> + <param name="idle_core_power" value="1.59"/><!-- idle core power for GTX479 --> + <!--param name="scaling_coefficients" value="10,0.0884816,10,10,8,10,4.12782,10,2.48832,10,10,10,4.29982,0.387764,0.0714269,0.14302,0.01,0.546811,0.485351,0.806633,0.818073,1.9207,100,100,100,87.9303,100,10,4.3548,10"/--> + <param name="TOT_INST" value="10" /> + <param name="FP_INT" value="10" /> + <param name="IC_H" value="0.001" /> + <param name="IC_M" value="10" /> + <param name="DC_RH" value="1" /> + <param name="DC_RM" value="1" /> + <param name="DC_WH" value="1" /> + <param name="DC_WM" value="1" /> + <param name="TC_H" value="0.001" /> + <param name="TC_M" value="10" /> + <param name="CC_H" value="4.5071" /> + <param name="CC_M" value="10" /> + <param name="SHRD_ACC" value="10" /> + <param name="REG_RD" value="1.6294" /> + <param name="REG_WR" value="0.5031" /> + <param name="NON_REG_OPs" value="0.01" /> + <param name="SP_ACC" value="10" /> + <param name="SFU_ACC" value="0.0082" /> + <param name="FPU_ACC" value="0.4126" /> + <param name="MEM_RD" value="0.1234" /> + <param name="MEM_WR" value="0.001" /> + <param name="MEM_PRE" value="0.001" /> + <param name="L2_RH" value="100" /> + <param name="L2_RM" value="100" /> + <param name="L2_WH" value="100" /> + <param name="L2_WM" value="42.6966" /> + <param name="NOC_A" value="100" /> + <param name="PIPE_A" value="44.8085" /> + <param name="IDLE_CORE_N" value="2.0382"/> + <param name="CONST_DYNAMICN" value="5.0005" /> + <stat name="num_idle_cores" value="0"/><!-- Average Number of idle cores during this period --> + <stat name="total_cycles" value="total_cycles_match_mcpat"/> + <stat name="idle_cycles" value="idle_cycles_match_mcpat"/> + <stat name="busy_cycles" value="busy_cycles_match_mcpat"/> + <!--This page size(B) is complete different from the page size in Main memo secction. this page size is the size of + virtual memory from OS/Archi perspective; the page size in Main memo secction is the actuall physical line in a DRAM bank --> + <!-- *********************** cores ******************* --> + <component id="system.core0" name="core0"> + <!-- Core property --> + <param name="clock_rate" value="1481"/> + <param name="instruction_length" value="32"/> + <param name="opcode_width" value="9"/> + <!-- address width determins the tag_width in Cache, LSQ and buffers in cache controller + default value is machine_bits, if not set --> + <param name="machine_type" value="1"/><!-- 1 inorder; 0 OOO--> + <!-- inorder/OoO --> + <param name="number_hardware_threads" value="32"/> + <!-- number_instruction_fetch_ports(icache ports) is always 1 in single-thread processor, + it only may be more than one in SMT processors. BTB ports always equals to fetch ports since + branch information in consective branch instructions in the same fetch group can be read out from BTB once.--> + <param name="fetch_width" value="1"/> + <!-- fetch_width determins the size of cachelines of L1 cache block --> + <param name="number_instruction_fetch_ports" value="1"/> + <param name="decode_width" value="1"/> + <!-- decode_width determins the number of ports of the + renaming table (both RAM and CAM) scheme --> + <param name="issue_width" value="2"/> + <!-- issue_width determins the number of ports of Issue window and other logic + as in the complexity effective proccessors paper; issue_width==dispatch_width --> + <param name="commit_width" value="2"/> + <!-- commit_width determins the number of ports of register files --> + <param name="fp_issue_width" value="1"/> + <param name="prediction_width" value="0"/> + <!-- number of branch instructions can be predicted simultannouesl--> + <!-- Current version of McPAT does not distinguish int and floating point pipelines + Theses parameters are reserved for future use.--> + <param name="pipelines_per_core" value="1,1"/> + <!--integer_pipeline and floating_pipelines, if the floating_pipelines is 0, then the pipeline is shared--> + <param name="pipeline_depth" value="8,8"/> + <!-- pipeline depth of int and fp, if pipeline is shared, the second number is the average cycles of fp ops --> + <!-- issue and exe unit--> + <param name="ALU_per_core" value="32"/> + <!-- contains an adder, a shifter, and a logical unit --> + <param name="MUL_per_core" value="4"/> + <!-- For MUL and Div --> + <param name="FPU_per_core" value="32"/> + <!-- buffer between IF and ID stage --> + <param name="instruction_buffer_size" value="1"/> + <!-- buffer between ID and sche/exe stage --> + <param name="decoded_stream_buffer_size" value="1"/> + <param name="instruction_window_scheme" value="0"/><!-- 0 PHYREG based, 1 RSBASED--> + <!-- McPAT support 2 types of OoO cores, RS based and physical reg based--> + <param name="instruction_window_size" value="1"/> + <param name="fp_instruction_window_size" value="1"/> + <!-- the instruction issue Q as in Alpha 21264; The RS as in Intel P6 --> + <param name="ROB_size" value="0"/> + <!-- each in-flight instruction has an entry in ROB --> + <!-- registers --> + <!-- SM parameters Added by Syed Gilani --> + <param name="rf_banks" value="32"/> + <param name="simd_width" value="32"/> + <param name="collector_units" value="32"/> + <param name="core_clock_ratio" value="2"/> + <param name="warp_size" value="32"/> + + <param name="archi_Regs_IRF_size" value="65536"/> + <param name="archi_Regs_FRF_size" value="32"/> + <!-- if OoO processor, phy_reg number is needed for renaming logic, + renaming logic is for both integer and floating point insts. --> + <param name="phy_Regs_IRF_size" value="32"/> + <param name="phy_Regs_FRF_size" value="32"/> + <!-- rename logic --> + <param name="rename_scheme" value="0"/> + <!-- can be RAM based(0) or CAM based(1) rename scheme + RAM-based scheme will have free list, status table; + CAM-based scheme have the valid bit in the data field of the CAM + both RAM and CAM need RAM-based checkpoint table, checkpoint_depth=# of in_flight instructions; + Detailed RAT Implementation see TR --> + <param name="register_windows_size" value="0"/> + <!-- how many windows in the windowed register file, sun processors; + no register windowing is used when this number is 0 --> + <!-- In OoO cores, loads and stores can be issued whether inorder(Pentium Pro) or (OoO)out-of-order(Alpha), + They will always try to exeute out-of-order though. --> + <param name="LSU_order" value="inorder"/> + <param name="store_buffer_size" value="32"/> + <!-- By default, in-order cores do not have load buffers --> + <param name="load_buffer_size" value="32"/> + <!-- number of ports refer to sustainable concurrent memory accesses --> + <param name="memory_ports" value="2"/> + <!-- max_allowed_in_flight_memo_instructions determins the # of ports of load and store buffer + as well as the ports of Dcache which is connected to LSU --> + <!-- dual-pumped Dcache can be used to save the extra read/write ports --> + <param name="RAS_size" value="1"/> + <!-- general stats, defines simulation periods;require total, idle, and busy cycles for senity check --> + <!-- please note: if target architecture is X86, then all the instrucions refer to (fused) micro-ops --> + <stat name="total_instructions" value="total_instructions_match_mcpat"/> + <stat name="int_instructions" value="int_instruction_match_mcpat"/> + <stat name="fp_instructions" value="flt_instruction_match_mcpat"/> + <stat name="branch_instructions" value="branch_instruction_match_mcpat"/> + <stat name="branch_mispredictions" value="0"/> + <stat name="load_instructions" value="load_instruction_match_mcpat"/> + <stat name="store_instructions" value="store_instruction_match_mcpat"/> + <stat name="committed_instructions" value="total_instructions_match_mcpat"/> + <stat name="committed_int_instructions" value="int_instruction_match_mcpat"/> + <stat name="committed_fp_instructions" value="flt_instruction_match_mcpat"/> + <stat name="pipeline_duty_cycle" value="0.6"/><!--<=1, runtime_ipc/peak_ipc; averaged for all cores if homogenous --> + <!-- the following cycle stats are used for heterogeneouse cores only, + please ignore them if homogeneouse cores --> + <stat name="total_cycles" value="total_cycles_match_mcpat"/> + <stat name="idle_cycles" value="idle_cycles_match_mcpat"/> + <stat name="busy_cycles" value="busy_cycles_match_mcpat"/> + <!-- instruction buffer stats --> + <!-- ROB stats, both RS and Phy based OoOs have ROB + performance simulator should capture the difference on accesses, + otherwise, McPAT has to guess based on number of commited instructions. --> + <stat name="ROB_reads" value="263886"/> + <stat name="ROB_writes" value="263886"/> + <!-- RAT accesses --> + <stat name="rename_accesses" value="263886"/> + <stat name="fp_rename_accesses" value="263886"/> + <!-- decode and rename stage use this, should be total ic - nop --> + <!-- Inst window stats --> + <stat name="inst_window_reads" value="263886"/> + <stat name="inst_window_writes" value="263886"/> + <stat name="inst_window_wakeup_accesses" value="263886"/> + <stat name="fp_inst_window_reads" value="263886"/> + <stat name="fp_inst_window_writes" value="263886"/> + <stat name="fp_inst_window_wakeup_accesses" value="263886"/> + <!-- RF accesses --> + <stat name="int_regfile_reads" value="int_register_read_access_match_mcpat"/> + <stat name="float_regfile_reads" value="int_register_write_access_match_mcpat"/> + <stat name="int_regfile_writes" value="float_register_read_access_match_mcpat"/> + <stat name="float_regfile_writes" value="float_register_write_access_match_mcpat"/> + + <!-- The following stat is for operand collector power - Added by Syed --> + <stat name="non_rf_operands" value="0"/> + + <!-- accesses to the working reg --> + <stat name="function_calls" value="0"/> + <stat name="context_switches" value="0"/> <!--not used in the McPAT --> + <!-- Number of Windowes switches (number of function calls and returns)--> + <!-- Alu stats by default, the processor has one FPU that includes the divider and + multiplier. The fpu accesses should include accesses to multiplier and divider --> + <stat name="ialu_accesses" value="ialu_accesses_match_mcpat"/> + <stat name="fpu_accesses" value="fpu_accesses_match_mcpat"/> + <stat name="mul_accesses" value="mul_accesses_match_mcpat"/> + <stat name="cdb_alu_accesses" value="0"/> + <stat name="cdb_mul_accesses" value="0"/> + <stat name="cdb_fpu_accesses" value="0"/> + <!-- multiple cycle accesses should be counted multiple times, + otherwise, McPAT can use internal counter for different floating point instructions + to get final accesses. But that needs detailed info for floating point inst mix --> + <!-- currently the performance simulator should + make sure all the numbers are final numbers, + including the explicit read/write accesses, + and the implicite accesses such as replacements and etc. + Future versions of McPAT may be able to reason the implicite access + based on param and stats of last level cache + The same rule applies to all cache access stats too! --> + <!-- following is AF for max power computation. + Do not change them, unless you understand them--> + <stat name="IFU_duty_cycle" value="0.25"/> + <stat name="LSU_duty_cycle" value="0.25"/> + <stat name="MemManU_I_duty_cycle" value="1"/> + <stat name="MemManU_D_duty_cycle" value="0.25"/> + <stat name="ALU_duty_cycle" value="0.9"/> + <stat name="MUL_duty_cycle" value="0.5"/> + <stat name="FPU_duty_cycle" value="1"/><!-- FPU numbers are already average --> + <stat name="ALU_cdb_duty_cycle" value="0.9"/> + <stat name="MUL_cdb_duty_cycle" value="0.5"/> + <stat name="FPU_cdb_duty_cycle" value="15"/> + <component id="system.core0.predictor" name="PBT"> + <!-- branch predictor; tournament predictor see Alpha implementation --> + <param name="local_predictor_size" value="10,3"/> + <param name="local_predictor_entries" value="1024"/> + <param name="global_predictor_entries" value="4096"/> + <param name="global_predictor_bits" value="2"/> + <param name="chooser_predictor_entries" value="4096"/> + <param name="chooser_predictor_bits" value="2"/> + <!-- These parameters can be combined like below in next version + <param name="load_predictor" value="10,3,1024"/> + <param name="global_predictor" value="4096,2"/> + <param name="predictor_chooser" value="4096,2"/> + --> + </component> + <component id="system.core0.itlb" name="itlb"> + <param name="number_entries" value="1"/> + <stat name="total_accesses" value="0"/> + <stat name="total_misses" value="0"/> + <stat name="conflicts" value="0"/> + <!-- there is no write requests to itlb although writes happen to itlb after miss, + which is actually a replacement --> + </component> + <component id="system.core0.icache" name="icache"> + <!-- there is no write requests to itlb although writes happen to it after miss, + which is actually a replacement --> + <param name="icache_config" value="16384,128,4,1,1,3,8,0"/> + <!-- the parameters are capacity,block_width, associativity, bank, throughput w.r.t. core clock, latency w.r.t. core clock,output_width, cache policy --> + <!-- cache_policy;//0 no write or write-though with non-write allocate;1 write-back with write-allocate --> + <param name="buffer_sizes" value="16, 16, 16,0"/> + <!-- cache controller buffer sizes: miss_buffer_size(MSHR),fill_buffer_size,prefetch_buffer_size,wb_buffer_size--> + <stat name="read_accesses" value="total_instructions_match_mcpat"/> + <stat name="read_misses" value="0"/> + <stat name="conflicts" value="0"/> + </component> + <component id="system.core0.dtlb" name="dtlb"> + <param name="number_entries" value="1"/> + <stat name="total_accesses" value="0"/> + <stat name="total_misses" value="0"/> + <stat name="conflicts" value="0"/> + </component> + <component id="system.core0.ccache" name="ccache"> + <!-- all the buffer related are optional --> + <param name="ccache_config" value="16384,64,2,1,1,3,8,0"/> + <param name="buffer_sizes" value="16, 16, 16, 0"/> + <!-- cache controller buffer sizes: miss_buffer_size(MSHR),fill_buffer_size,prefetch_buffer_size,wb_buffer_size--> + <stat name="read_accesses" value="ccache_read_accesses_match_mcpat"/> + <stat name="write_accesses" value="0"/> + <stat name="read_misses" value="ccache_read_misses_match_mcpat"/> + <stat name="write_misses" value="0"/> + <stat name="conflicts" value="0"/> + </component> + <component id="system.core0.tcache" name="tcache"> + <!-- all the buffer related are optional --> + <param name="tcache_config" value="49152,128,8,1,1,3,8,0"/> + <param name="buffer_sizes" value="16, 16, 16, 0"/> + <!-- cache controller buffer sizes: miss_buffer_size(MSHR),fill_buffer_size,prefetch_buffer_size,wb_buffer_size--> + <stat name="read_accesses" value="tcache_read_accesses_match_mcpat"/> + <stat name="write_accesses" value="0"/> + <stat name="read_misses" value="tcache_read_misses_match_mcpat"/> + <stat name="write_misses" value="0"/> + <stat name="conflicts" value="0"/> + </component> + <!--model the shared memory by mimicing dcache--> + <component id="system.core0.sharedmemory" name="sharedmemory"> + <!-- all the buffer related are optional --> + <param name="sharedmemory_config" value="98304,16,1,16,1,3,16,0"/> + <!-- the parameters are capacity,block_width, associativity, bank, throughput w.r.t. core clock, latency w.r.t. core clock,output_width, cache policy --> + <param name="buffer_sizes" value="16, 16, 16, 16"/> + <!-- cache controller buffer sizes: miss_buffer_size(MSHR),fill_buffer_size,prefetch_buffer_size,wb_buffer_size--> + <stat name="read_accesses" value="sharedmemory_read_access_match_mcpat"/> + <stat name="write_accesses" value="sharedmemory_write_access_match_mcpat"/> + <stat name="read_misses" value="0"/> + <stat name="write_misses" value="0"/> + <stat name="conflicts" value="0"/> + </component> + <component id="system.core0.dcache" name="dcache"> + <!-- all the buffer related are optional --> + <param name="dcache_config" value="16384,32,4,1,1,3,8,0"/> + <param name="buffer_sizes" value="16, 16, 16, 0"/> + <!-- cache controller buffer sizes: miss_buffer_size(MSHR),fill_buffer_size,prefetch_buffer_size,wb_buffer_size--> + <stat name="read_accesses" value="dcache_read_access_match_mcpat"/> + <stat name="write_accesses" value="dcache_write_access_match_mcpat"/> + <stat name="read_misses" value="dcache_read_miss_match_mcpat"/> + <stat name="write_misses" value="dcache_write_miss_match_mcpat"/> + <stat name="conflicts" value="0"/> + </component> + <component id="system.core0.BTB" name="BTB"> + <!-- all the buffer related are optional --> + <param name="BTB_config" value="8192,4,2,1, 1,3"/> + <!-- the parameters are capacity,block_width,associativity,bank, throughput w.r.t. core clock, latency w.r.t. core clock,--> + </component> + </component> + <component id="system.L1Directory0" name="L1Directory0"> + <param name="Directory_type" value="0"/> + <!--0 cam based shadowed tag. 1 directory cache --> + <param name="Dir_config" value="2048,1,0,1, 4, 4,8"/> + <!-- the parameters are capacity,block_width, associativity,bank, throughput w.r.t. core clock, latency w.r.t. core clock,--> + <param name="buffer_sizes" value="8, 8, 8, 8"/> + <!-- all the buffer related are optional --> + <param name="clockrate" value="1400"/> + <param name="ports" value="1,1,1"/> + <!-- number of r, w, and rw search ports --> + <param name="device_type" value="0"/> + <!-- altough there are multiple access types, + Performance simulator needs to cast them into reads or writes + e.g. the invalidates can be considered as writes --> + <stat name="read_accesses" value="800000"/> + <stat name="write_accesses" value="27276"/> + <stat name="read_misses" value="1632"/> + <stat name="write_misses" value="183"/> + <stat name="conflicts" value="20"/> + <stat name="duty_cycle" value="0.45"/> + </component> + <component id="system.L2Directory0" name="L2Directory0"> + <param name="Directory_type" value="1"/> + <!--0 cam based shadowed tag. 1 directory cache --> + <param name="Dir_config" value="1048576,16,16,1,2, 100"/> + <!-- the parameters are capacity,block_width, associativity,bank, throughput w.r.t. core clock, latency w.r.t. core clock,--> + <param name="buffer_sizes" value="8, 8, 8, 8"/> + <!-- all the buffer related are optional --> + <param name="clockrate" value="1400"/> + <param name="ports" value="1,1,1"/> + <!-- number of r, w, and rw search ports --> + <param name="device_type" value="0"/> + <!-- altough there are multiple access types, + Performance simulator needs to cast them into reads or writes + e.g. the invalidates can be considered as writes --> + <stat name="read_accesses" value="0"/> + <stat name="write_accesses" value="0"/> + <stat name="read_misses" value="0"/> + <stat name="write_misses" value="0"/> + <stat name="conflicts" value="0"/> + <stat name="duty_cycle" value="0.45"/> + </component> + <component id="system.L20" name="L20"> + <!-- all the buffer related are optional --> + <param name="L2_config" value="131072,128,16,1, 4,23, 64, 1"/> + <!-- consider 4-way bank interleaving for Niagara 1 --> + <!-- the parameters are capacity,block_width, associativity, bank, throughput w.r.t. core clock, latency w.r.t. core clock,output_width, cache policy --> + <param name="buffer_sizes" value="16, 16, 16, 16"/> + <!-- cache controller buffer sizes: miss_buffer_size(MSHR),fill_buffer_size,prefetch_buffer_size,wb_buffer_size--> + <param name="clockrate" value="2962"/> + <param name="ports" value="1,1,1"/> + <!-- number of r, w, and rw ports --> + <param name="device_type" value="0"/> + <stat name="read_accesses" value="200000"/> + <stat name="write_accesses" value="0"/> + <stat name="read_misses" value="0"/> + <stat name="write_misses" value="0"/> + <stat name="conflicts" value="0"/> + <stat name="duty_cycle" value="0.5"/> + </component> + +<!--**********************************************************************--> +<component id="system.L30" name="L30"> + <param name="L3_config" value="1048576,64,16,1, 2,100, 64,1"/> + <!-- the parameters are capacity,block_width, associativity, bank, throughput w.r.t. core clock, latency w.r.t. core clock,output_width, cache policy --> + <param name="clockrate" value="3500"/> + <param name="ports" value="1,1,1"/> + <!-- number of r, w, and rw ports --> + <param name="device_type" value="0"/> + <param name="buffer_sizes" value="16, 16, 16, 16"/> + <!-- cache controller buffer sizes: miss_buffer_size(MSHR),fill_buffer_size,prefetch_buffer_size,wb_buffer_size--> + <stat name="read_accesses" value="58824"/> + <stat name="write_accesses" value="27276"/> + <stat name="read_misses" value="1632"/> + <stat name="write_misses" value="183"/> + <stat name="conflicts" value="0"/> + <stat name="duty_cycle" value="0.35"/> + </component> + + +<!--**********************************************************************--> + <component id="system.NoC0" name="noc0"> + <param name="clockrate" value="700"/> + <param name="type" value="1"/> + <!-- 1 NoC, O bus --> + <param name="horizontal_nodes" value="2"/> + <param name="vertical_nodes" value="1"/> + <param name="has_global_link" value="0"/> + <!-- 1 has global link, 0 does not have global link --> + <param name="link_throughput" value="1"/><!--w.r.t clock --> + <param name="link_latency" value="1"/><!--w.r.t clock --> + <!-- througput >= latency --> + <!-- Router architecture --> + <param name="input_ports" value="6"/> + <param name="output_ports" value="6"/> + <param name="virtual_channel_per_port" value="1"/> + <!-- input buffer; in classic routers only input ports need buffers --> + <param name="flit_bits" value="32"/> + <param name="input_buffer_entries_per_vc" value="1"/><!--VCs within the same ports share input buffers whose size is propotional to the number of VCs--> + <param name="chip_coverage" value="1"/> + <!-- When multiple NOC present, one NOC will cover part of the whole chip. chip_coverage <=1 --> + <stat name="total_accesses" value="0"/> + <!-- This is the number of total accesses within the whole network not for each router --> + <stat name="duty_cycle" value="0.6"/> + </component> +<!--**********************************************************************--> +<!--**********************************************************************--> + + <component id="system.mem" name="mem"> + <!-- Main memory property --> + <param name="mem_tech_node" value="23"/> + <param name="device_clock" value="200"/><!--MHz, this is clock rate of the actual memory device, not the FSB --> + <param name="peak_transfer_rate" value="3200"/><!--MB/S--> + <param name="internal_prefetch_of_DRAM_chip" value="4"/> + <!-- 2 for DDR, 4 for DDR2, 8 for DDR3...--> + <!-- the device clock, peak_transfer_rate, and the internal prefetch decide the DIMM property --> + <!-- above numbers can be easily found from Wikipedia --> + <param name="capacity_per_channel" value="4096"/> <!-- MB --> + <!-- capacity_per_Dram_chip=capacity_per_channel/number_of_dimms/number_ranks/Dram_chips_per_rank + Current McPAT assumes single DIMMs are used.--> + <param name="number_ranks" value="2"/> + <param name="num_banks_of_DRAM_chip" value="6"/> + <param name="Block_width_of_DRAM_chip" value="64"/> <!-- B --> + <param name="output_width_of_DRAM_chip" value="8"/> + <!--number of Dram_chips_per_rank=" 72/output_width_of_DRAM_chip--> + <!--number of Dram_chips_per_rank=" 72/output_width_of_DRAM_chip--> + <param name="page_size_of_DRAM_chip" value="8"/> <!-- 8 or 16 --> + <param name="burstlength_of_DRAM_chip" value="8"/> + <stat name="memory_accesses" value="1052"/> + <stat name="memory_reads" value="1052"/> + <stat name="memory_writes" value="1052"/> + </component> + <component id="system.mc" name="mc"> + <!-- Memeory controllers are for DDR(2,3...) DIMMs --> + <!-- current version of McPAT uses published values for base parameters of memory controller + improvments on MC will be added in later versions. --> + <param name="type" value="0"/> <!-- 1: low power; 0 high performance --> + <param name="mc_clock" value="1848"/><!--DIMM IO bus clock rate MHz DDR2-400 for Niagara 1--> + <param name="peak_transfer_rate" value="29568"/><!--MB/S Syed: GTX 470 has 177.4GB/s mem transfer rate with 6 MCs --> + <param name="block_size" value="64"/><!--B--> + <param name="number_mcs" value="6"/><!-- 6 GDDR5 memory controllers --> + <!-- current McPAT only supports homogeneous memory controllers --> + <param name="memory_channels_per_mc" value="2"/> + <param name="number_ranks" value="1"/> + <param name="withPHY" value="0"/> + <!-- # of ranks of each channel--> + <param name="req_window_size_per_channel" value="16"/> + <param name="IO_buffer_size_per_channel" value="16"/> + <param name="databus_width" value="32"/> + <param name="addressbus_width" value="32"/> + <param name="PRT_entries" value="32"/> + <!-- # of empirical DRAM model parameter --> + <param name="dram_cmd_coeff" value="0"/> + <param name="dram_act_coeff" value="0"/> + <param name="dram_nop_coeff" value="0"/> + <param name="dram_activity_coeff" value="0"/> + <param name="dram_pre_coeff" value="3.8475e-8f"/> + <param name="dram_rd_coeff" value="7.74707143e-8f"/> + <param name="dram_wr_coeff" value="3.54664286e-8f"/> + <param name="dram_req_coeff" value="0"/> + <param name="dram_const_coeff" value="0"/> + + <!-- McPAT will add the control bus width to the addressbus width automatically --> + <stat name="memory_accesses" value="memory_accesses_match_mcpat"/> + <stat name="memory_reads" value="memory_reads_match_mcpat"/> + <stat name="memory_writes" value="memory_writes_match_mcpat"/> + <!-- McPAT does not track individual mc, instead, it takes the total accesses and calculate + the average power per MC or per channel. This is sufficent for most application. + Further trackdown can be easily added in later versions. --> + </component> +<!--**********************************************************************--> + <component id="system.niu" name="niu"> + <!-- On chip 10Gb Ethernet NIC, including XAUI Phy and MAC controller --> + <!-- For a minimum IP packet size of 84B at 10Gb/s, a new packet arrives every 67.2ns. + the low bound of clock rate of a 10Gb MAC is 150Mhz --> + <param name="type" value="0"/> <!-- 1: low power; 0 high performance --> + <param name="clockrate" value="350"/> + <param name="number_units" value="0"/> <!-- unlike PCIe and memory controllers, each Ethernet controller only have one port --> + <stat name="duty_cycle" value="1.0"/> <!-- achievable max load <= 1.0 --> + <stat name="total_load_perc" value="0.7"/> <!-- ratio of total achived load to total achivable bandwidth --> + <!-- McPAT does not track individual nic, instead, it takes the total accesses and calculate + the average power per nic or per channel. This is sufficent for most application. --> + </component> +<!--**********************************************************************--> + <component id="system.pcie" name="pcie"> + <!-- On chip PCIe controller, including Phy--> + <!-- For a minimum PCIe packet size of 84B at 8Gb/s per lane (PCIe 3.0), a new packet arrives every 84ns. + the low bound of clock rate of a PCIe per lane logic is 120Mhz --> + <param name="type" value="0"/> <!-- 1: low power; 0 high performance --> + <param name="withPHY" value="1"/> + <param name="clockrate" value="350"/> + <param name="number_units" value="0"/> + <param name="num_channels" value="8"/> <!-- 2 ,4 ,8 ,16 ,32 --> + <stat name="duty_cycle" value="1.0"/> <!-- achievable max load <= 1.0 --> + <stat name="total_load_perc" value="0.7"/> <!-- Percentage of total achived load to total achivable bandwidth --> + <!-- McPAT does not track individual pcie controllers, instead, it takes the total accesses and calculate + the average power per pcie controller or per channel. This is sufficent for most application. --> + </component> +<!--**********************************************************************--> + <component id="system.flashc" name="flashc"> + <param name="number_flashcs" value="0"/> + <param name="type" value="1"/> <!-- 1: low power; 0 high performance --> + <param name="withPHY" value="1"/> + <param name="peak_transfer_rate" value="200"/><!--Per controller sustainable reak rate MB/S --> + <stat name="duty_cycle" value="1.0"/> <!-- achievable max load <= 1.0 --> + <stat name="total_load_perc" value="0.7"/> <!-- Percentage of total achived load to total achivable bandwidth --> + <!-- McPAT does not track individual flash controller, instead, it takes the total accesses and calculate + the average power per fc or per channel. This is sufficent for most application --> + </component> +<!--**********************************************************************--> + + </component> +</component> diff --git a/cuda-kernels/tensor_core b/cuda-kernels/tensor_core Binary files differnew file mode 100755 index 0000000..b25f3d9 --- /dev/null +++ b/cuda-kernels/tensor_core diff --git a/cuda-kernels/tensor_core.cu b/cuda-kernels/tensor_core.cu new file mode 100644 index 0000000..483a42b --- /dev/null +++ b/cuda-kernels/tensor_core.cu @@ -0,0 +1,250 @@ +/* Copyright (c) 1993-2017, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +#include <stdio.h> + +// Define some error checking macros. +#define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } +void cudaErrCheck_(cudaError_t stat, const char *file, int line) { + if (stat != cudaSuccess) { + fprintf(stderr, "CUDA Error: %s %s %d\n", cudaGetErrorString(stat), file, line); + } +} + + + + +#include <mma.h> +using namespace nvcuda; + +// Must be multiples of 16 for wmma code to work +#define MATRIX_M (16) +#define MATRIX_N (16) +#define MATRIX_K (16) + + + +// The only dimensions currently supported by WMMA +const int WMMA_M = 16; +const int WMMA_N = 16; +const int WMMA_K = 16; + + +// Performs an MxNxK GEMM (C=alpha*A*B + beta*C) assuming: +// 1) Matrices are packed in memory. +// 2) M, N and K are multiples of 16. +// 3) Neither A nor B are transposed. +// Note: This is NOT a high performance example but is for demonstration purposes only +// For a high performance code please use the GEMM provided in cuBLAS. +__global__ void wmma_example(half *a, half *b, float *c, int M, int N, int K, float alpha, float beta) { + unsigned int start_time=0,end_time=0; + // Leading dimensions. Packed with no transpositions. + start_time=clock(); + int lda = M; + int ldb = K; + int ldc = M; + + // Tile using a 2D grid/ + int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / warpSize; + int warpN = (blockIdx.y * blockDim.y + threadIdx.y); + + // Declare the fragments + wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, half, wmma::row_major> a_frag; + wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, half, wmma::col_major> b_frag; + wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> acc_frag; + wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag; + + wmma::fill_fragment(c_frag, 0.0f); + + int i=0; + int aRow = warpM * WMMA_M; + int bCol = warpN * WMMA_N; + int aCol = i; + int bRow = i; + + + // Bounds checking + if (aRow < M && aCol < K && bRow < K && bCol < N) { + wmma::load_matrix_sync(a_frag, a+aRow+aCol*lda, lda); + wmma::load_matrix_sync(b_frag, b+bRow*ldb+bCol, ldb); + wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); + //wmma::mma_sync(acc_frag, a_frag, b_frag, acc_frag); + } + int cRow = warpM * WMMA_M; + int cCol = warpN * WMMA_N; + wmma::store_matrix_sync(c + cRow + cCol * ldc, c_frag, ldc, wmma::mem_col_major); + end_time=clock(); + printf("clock=%d",end_time-start_time); +} + +__global__ void convertFp32ToFp16 (half *out, float *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n) { + out[idx] = in[idx]; + } +} + +int main(int argc, char* argv[]) { + float *a_fp32; + float *b_fp32; + half *a_fp16; + half *b_fp16; + + float *c; + float *c_cublas; + float *c_wmma; + + float *c_host_cublas; + float *c_host_wmma; + float *a_host_wmma; + float *b_host_wmma; + float *c_init_host_wmma; + + + cudaEvent_t startWMMA; + cudaEvent_t stopWMMA; + + + cudaErrCheck(cudaEventCreate(&startWMMA)); + cudaErrCheck(cudaEventCreate(&stopWMMA)); + + + + + // Use tensor cores + + + cudaErrCheck(cudaMalloc((void**)&a_fp32, MATRIX_M * MATRIX_K * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&b_fp32, MATRIX_K * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&a_fp16, MATRIX_M * MATRIX_K * sizeof(half))); + cudaErrCheck(cudaMalloc((void**)&b_fp16, MATRIX_K * MATRIX_N * sizeof(half))); + + cudaErrCheck(cudaMalloc((void**)&c, MATRIX_M * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&c_wmma, MATRIX_M * MATRIX_N * sizeof(float))); + + c_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + c_init_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + a_host_wmma = (float*)malloc(MATRIX_M * MATRIX_K * sizeof(float)); + b_host_wmma = (float*)malloc(MATRIX_K * MATRIX_N * sizeof(float)); + + + +// printf("a_fp32\n"); + for(int m=0;m<MATRIX_M;m++){ + for(int n=0;n<MATRIX_K;n++){ + a_host_wmma[m*MATRIX_K+n]=m*MATRIX_K+n; + } + //printf(";\n"); + } + // printf("b_fp32\n"); + for(int m=0;m<MATRIX_K;m++){ + for(int n=0;n<MATRIX_N;n++){ + b_host_wmma[m*MATRIX_N+n]=m*MATRIX_N+n; +// printf("%f ",b_host_wmma[m*MATRIX_N+n]); + } +// printf(";\n"); + } + cudaErrCheck(cudaMemcpy(a_fp32,a_host_wmma, MATRIX_M * MATRIX_K * sizeof(float), cudaMemcpyHostToDevice)); + cudaErrCheck(cudaMemcpy(b_fp32,b_host_wmma, MATRIX_K * MATRIX_N * sizeof(float), cudaMemcpyHostToDevice)); + + // curand doesn't currently support fp16 so we generate in fp32 and convert to fp16. + convertFp32ToFp16 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_fp16, a_fp32, MATRIX_M * MATRIX_K); + convertFp32ToFp16 <<< (MATRIX_K * MATRIX_N + 255) / 256, 256 >>> (b_fp16, b_fp32, MATRIX_K * MATRIX_N); + + for(int m=0;m<MATRIX_M;m++){ + for(int n=0;n<MATRIX_N;n++){ + c_init_host_wmma[m*MATRIX_N+n]=0; + } + } + cudaErrCheck(cudaMemcpy(c, c_init_host_wmma, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyHostToDevice)); + cudaErrCheck(cudaMemcpy(c_wmma, c, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToDevice)); + + float alpha = 1.0f; + float beta = 1.0f; + + + printf("\nM = %d, N = %d, K = %d. alpha = %f, beta = %f\n\n", MATRIX_M, MATRIX_N, MATRIX_K, alpha, beta); + + // First: using WMMA + dim3 gridDim; + dim3 blockDim; + + // blockDim.x must be a multple of warpSize + // 128x4 means we have 16 warps and a block computes a 64x64 output tile + blockDim.x = 128; + blockDim.y = 4; + + gridDim.x = (MATRIX_M + (WMMA_M * blockDim.x / 32 - 1)) / (WMMA_M * blockDim.x / 32); + gridDim.y = (MATRIX_N + WMMA_N * blockDim.y - 1) / (WMMA_N * blockDim.y); + + printf("Running with wmma...\n"); + cudaErrCheck(cudaEventRecord(startWMMA)); + wmma_example <<< 1, 32>>> (a_fp16, b_fp16, c_wmma, MATRIX_M, MATRIX_N, MATRIX_K, alpha, beta); + // wmma_example <<< gridDim, blockDim >>> (a_fp16, b_fp16, c_wmma, MATRIX_M, MATRIX_N, MATRIX_K, alpha, beta); + cudaErrCheck(cudaEventRecord(stopWMMA)); + + + + + // Error checking + printf("\nChecking results...\n"); + cudaErrCheck(cudaMemcpy(c_host_wmma, c_wmma, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost)); + // printf("c_host\n"); + // for(int m=0;m<MATRIX_M;m++){ +// for(int n=0;n<MATRIX_N;n++){ +// printf("%f ",c_host_wmma[m*MATRIX_N+n]); +// } +// printf(";\n"); + // } + + float wmmaTime; + cudaErrCheck(cudaEventSynchronize(stopWMMA)); + cudaErrCheck(cudaEventElapsedTime(&wmmaTime, startWMMA, stopWMMA)); + printf("wmma took %fms\n", wmmaTime); + //printf("Clock=%d",stopWMMA-startWMMA); + printf("\nFor a faster code using wmma you should check out the cudaTensorCoreGemm sample in the CUDA Toolkit.\nThis code was written as a demo only!\n\n"); + + + cudaErrCheck(cudaEventDestroy(startWMMA)); + cudaErrCheck(cudaEventDestroy(stopWMMA)); + + + cudaErrCheck(cudaFree(a_fp32)); + cudaErrCheck(cudaFree(b_fp32)); + cudaErrCheck(cudaFree(a_fp16)); + cudaErrCheck(cudaFree(b_fp16)); + + cudaErrCheck(cudaFree(c)); + cudaErrCheck(cudaFree(c_wmma)); + + free(c_host_wmma); + + cudaErrCheck(cudaDeviceReset()); + return 0; +} + + diff --git a/cuobjdump_to_ptxplus/ptx_parser.h b/cuobjdump_to_ptxplus/ptx_parser.h index 7cfb5c5..a534e92 100644 --- a/cuobjdump_to_ptxplus/ptx_parser.h +++ b/cuobjdump_to_ptxplus/ptx_parser.h @@ -82,6 +82,7 @@ void add_file( unsigned a, const char *b ) {PTX_PARSE_DPRINTF(" ");} void add_variables() {PTX_PARSE_DPRINTF(" ");} void set_variable_type() {PTX_PARSE_DPRINTF(" ");} void add_option(int a ) {PTX_PARSE_DPRINTF(" ");} +void add_wmma_option(int a ) {PTX_PARSE_DPRINTF(" ");} void add_array_initializer() {PTX_PARSE_DPRINTF(" ");} void add_label( const char *a ) {PTX_PARSE_DPRINTF(" ");} void set_return() {PTX_PARSE_DPRINTF(" ");} @@ -99,6 +100,7 @@ void add_1vector_operand( const char *a ) {PTX_PARSE_DPRINTF(" ");} void add_2vector_operand( const char *a, const char *b ) {PTX_PARSE_DPRINTF(" ");} void add_3vector_operand( const char *a, const char *b, const char *c ) {PTX_PARSE_DPRINTF(" ");} void add_4vector_operand( const char *a, const char *b, const char *c, const char *d ) {PTX_PARSE_DPRINTF(" ");} +void add_8vector_operand( const char *a, const char *b, const char *c, const char *d ,const char *e,const char *f,const char *g,const char *h) {PTX_PARSE_DPRINTF(" ");} void add_builtin_operand( int a, int b ) {PTX_PARSE_DPRINTF(" ");} void add_memory_operand() {PTX_PARSE_DPRINTF(" ");} void change_memory_addr_space( const char *a ) {PTX_PARSE_DPRINTF(" ");} diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index bddaf8c..7407269 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1493,6 +1493,12 @@ unsigned trunc(unsigned num, unsigned precision) { } return num; } +void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +{ +} +void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +{ +} void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { @@ -1565,7 +1571,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) for (i=0;i<16;i++){ for(j=0;j<16;j++){ for(k=0;k<16;k++){ - matrix_d[i][j].f32=matrix_d[i][j].f32+matrix_a[i][k].f32*matrix_b[k][j].f32; + matrix_d[i][j].f32=matrix_d[i][j].f32+matrix_a[i][k].f32*matrix_b[j][k].f32; } matrix_d[i][j].f32+=matrix_c[i][j].f32; } diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index a3cc83f..e6f957a 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -53,6 +53,8 @@ OP_DEF(BRX_OP,brx_impl,"brx",0,3) OP_DEF(BREV_OP,brev_impl,"brev",1,1) OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9) OP_W_DEF(MMA_OP,mma_impl,"mma",1,1) +OP_W_DEF(MMA_LD_OP,mma_ld_impl,"mma_load",1,5) +OP_W_DEF(MMA_ST_OP,mma_st_impl,"mma_store",0,5) OP_DEF(CALL_OP,call_impl,"call",1,3) OP_DEF(CALLP_OP,callp_impl,"callp",1,3) OP_DEF(CLZ_OP,clz_impl,"clz",1,1) diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index aa133da..b91d92f 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -60,5 +60,14 @@ enum special_regs { WARPID_REG, WARPSZ_REG }; - +enum wmma_type{ + LOAD_A, + LOAD_B, + LOAD_C, + STORE_D, + MMA, + ROW, + COL, + M16N16K16 +}; #endif diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index e07e339..f67dddd 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -60,7 +60,7 @@ addc TC; ptx_lval.int_value = ADDC_OP; return OPCODE; and TC; ptx_lval.int_value = AND_OP; return OPCODE; andn TC; ptx_lval.int_value = ANDN_OP; return OPCODE; atom TC; ptx_lval.int_value = ATOM_OP; return OPCODE; -bar TC; ptx_lval.int_value = BAR_OP; return OPCODE; +bar TC; ptx_lval.int_value = BAR_OP; return OPCODE; bfe TC; ptx_lval.int_value = BFE_OP; return OPCODE; bfi TC; ptx_lval.int_value = BFI_OP; return OPCODE; bfind TC; ptx_lval.int_value = BFIND_OP; return OPCODE; @@ -68,9 +68,12 @@ bra TC; ptx_lval.int_value = BRA_OP; return OPCODE; brx TC; ptx_lval.int_value = BRX_OP; return OPCODE; brev TC; ptx_lval.int_value = BREV_OP; return OPCODE; brkpt TC; ptx_lval.int_value = BRKPT_OP; return OPCODE; -mma TC; ptx_lval.int_value = MMA_OP; return OPCODE; +wmma\.mma TC; ptx_lval.int_value = MMA_OP; return OPCODE; +wmma\.load TC; ptx_lval.int_value = MMA_LD_OP; return OPCODE; +wmma\.store TC; ptx_lval.int_value = MMA_ST_OP; return OPCODE; + call TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALL_OP; return OPCODE; // blocking opcode token in case the callee has the same name as an opcode -callp TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALLP_OP; return OPCODE; +callp TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALLP_OP; return OPCODE; clz TC; ptx_lval.int_value = CLZ_OP; return OPCODE; cnot TC; ptx_lval.int_value = CNOT_OP; return OPCODE; cos TC; ptx_lval.int_value = COS_OP; return OPCODE; @@ -153,6 +156,15 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; "CPTX_END" printf("ENDING CUSTOM PTX.\n"); BEGIN(IN_COMMENT); <INITIAL,NOT_OPCODE,IN_INST,IN_FUNC_DECL>{ +\.a\.sync TC; ptx_lval.int_value = LOAD_A; return WMMA_DIRECTIVE; +\.b\.sync TC; ptx_lval.int_value = LOAD_B; return WMMA_DIRECTIVE; +\.c\.sync TC; ptx_lval.int_value = LOAD_C; return WMMA_DIRECTIVE; +\.d\.sync TC; ptx_lval.int_value = STORE_D; return WMMA_DIRECTIVE; +\.sync TC;ptx_lval.int_value=MMA; return WMMA_DIRECTIVE; +\.row TC; ptx_lval.int_value = ROW; return LAYOUT; +\.col TC; ptx_lval.int_value = COL; return LAYOUT; +\.m16n16k16 TC; ptx_lval.int_value = M16N16K16; return CONFIGURATION; + \.align TC; return ALIGN_DIRECTIVE; \.branchtargets TC; return BRANCHTARGETS_DIRECTIVE; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 3360c55..737657c 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -37,6 +37,9 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token <string_value> STRING %token <int_value> OPCODE +%token <int_value> WMMA_DIRECTIVE +%token <int_value> LAYOUT +%token <int_value> CONFIGURATION %token ALIGN_DIRECTIVE %token BRANCHTARGETS_DIRECTIVE %token BYTE_DIRECTIVE @@ -428,6 +431,7 @@ option: type_spec | compare_spec | addressable_spec | rounding_mode + | wmma_spec | SYNC_OPTION { add_option(SYNC_OPTION); } | ARRIVE_OPTION { add_option(ARRIVE_OPTION); } | RED_OPTION { add_option(RED_OPTION); } @@ -483,6 +487,7 @@ atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } rounding_mode: floating_point_rounding_mode | integer_rounding_mode; + floating_point_rounding_mode: RN_OPTION { add_option(RN_OPTION); } | RZ_OPTION { add_option(RZ_OPTION); } | RM_OPTION { add_option(RM_OPTION); } @@ -515,6 +520,10 @@ compare_spec:EQ_OPTION { add_option(EQ_OPTION); } | NAN_OPTION { add_option(NAN_OPTION); } ; +wmma_spec: WMMA_DIRECTIVE LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);} + | WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2),add_wmma_option($3),add_wmma_option($4)} + ; + operand_list: operand | operand COMMA operand_list; @@ -543,6 +552,7 @@ operand: IDENTIFIER { add_scalar_operand( $1 ); } vector_operand: LEFT_BRACE IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_2vector_operand($2,$4); } | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_3vector_operand($2,$4,$6); } | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_4vector_operand($2,$4,$6,$8); } + | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_8vector_operand($2,$4,$6,$8,$10,$12,$14,$16); } | LEFT_BRACE IDENTIFIER RIGHT_BRACE { add_1vector_operand($2); } ; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 8ebdcf8..9a4d8d3 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -995,7 +995,7 @@ static std::list<operand_info> check_operands( int opcode, const std::list<operand_info> &operands ) { static int g_warn_literal_operands_two_type_inst; - if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) ) { + if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) || (opcode==MMA_OP)) { // just make sure these do not have have const operands... if( !g_warn_literal_operands_two_type_inst ) { std::list<operand_info>::const_iterator o; @@ -1043,6 +1043,7 @@ ptx_instruction::ptx_instruction( int opcode, const std::list<operand_info> &operands, const operand_info &return_var, const std::list<int> &options, + const std::list<int> &wmma_options, const std::list<int> &scalar_type, memory_space_t space_spec, const char *file, @@ -1061,6 +1062,7 @@ ptx_instruction::ptx_instruction( int opcode, m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() ); m_return_var = return_var; m_options = options; + m_wmma_options = wmma_options; m_wide = false; m_hi = false; m_lo = false; @@ -1078,7 +1080,7 @@ ptx_instruction::ptx_instruction( int opcode, m_atomic_spec = 0; m_membar_level = 0; m_inst_size = 8; // bytes - + int rr=0; std::list<int>::const_iterator i; unsigned n=1; for ( i=options.begin(); i!= options.end(); i++, n++ ) { diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 0601b97..ff24a66 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -582,6 +582,34 @@ public: m_is_return_var = false; m_immediate_address=false; } + operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 ,const symbol *s5,const symbol *s6,const symbol *s7, const symbol *s8) + { + init(); + m_is_non_arch_reg = false; + m_addr_space = undefined_space; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; + m_uid = get_uid(); + m_valid = true; + m_vector = true; + m_type = vector_t; + m_value.m_vector_symbolic = new const symbol*[8]; + m_value.m_vector_symbolic[0] = s1; + m_value.m_vector_symbolic[1] = s2; + m_value.m_vector_symbolic[2] = s3; + m_value.m_vector_symbolic[3] = s4; + m_value.m_vector_symbolic[4] = s5; + m_value.m_vector_symbolic[5] = s6; + m_value.m_vector_symbolic[6] = s7; + m_value.m_vector_symbolic[7] = s8; + m_addr_offset = 0; + m_neg_pred = false; + m_is_return_var = false; + m_immediate_address=false; + } + void init() { m_uid=(unsigned)-1; @@ -866,6 +894,7 @@ public: const std::list<operand_info> &operands, const operand_info &return_var, const std::list<int> &options, + const std::list<int> &wmma_options, const std::list<int> &scalar_type, memory_space_t space_spec, const char *file, @@ -1087,6 +1116,7 @@ private: operand_info m_return_var; std::list<int> m_options; + std::list<int> m_wmma_options; bool m_wide; bool m_hi; bool m_lo; @@ -1096,6 +1126,9 @@ private: bool m_uni; //if branch instruction, this evaluates to true for uniform branches (ie jumps) bool m_to_option; unsigned m_cache_option; + unsigned m_wmma_type; + unsigned m_wmma_layout[2]; + unsigned m_wmma_configuration; unsigned m_rounding_mode; unsigned m_compare_op; unsigned m_saturation_mode; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 7fc54e9..6757091 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -72,6 +72,7 @@ symbol *g_label; int g_opcode = -1; std::list<operand_info> g_operands; std::list<int> g_options; +std::list<int> g_wmma_options; std::list<int> g_scalar_type; #define PTX_PARSE_DPRINTF(...) \ @@ -162,6 +163,7 @@ void init_instruction_state() g_label = NULL; g_opcode = -1; g_options.clear(); + g_wmma_options.clear(); g_return_var = operand_info(); init_directive_state(); } @@ -300,6 +302,7 @@ void add_instruction() g_operands, g_return_var, g_options, + g_wmma_options, g_scalar_type, g_space_spec, g_filename, @@ -629,7 +632,7 @@ void add_scalar_type_spec( int type_spec ) g_scalar_type.push_back( type_spec ); if ( g_scalar_type.size() > 1 ) { parse_assert( (g_opcode == -1) || (g_opcode == CVT_OP) || (g_opcode == SET_OP) || (g_opcode == SLCT_OP) - || (g_opcode == TEX_OP), + || (g_opcode == TEX_OP)|| (g_opcode==MMA_OP), "only cvt, set, slct, and tex can have more than one type specifier."); } g_scalar_type_spec = type_spec; @@ -669,7 +672,11 @@ void add_option( int option ) PTX_PARSE_DPRINTF("add_option"); g_options.push_back( option ); } - +void add_wmma_option( int option ) +{ + PTX_PARSE_DPRINTF("add_option"); + g_wmma_options.push_back( option ); +} void add_double_operand( const char *d1, const char *d2 ) { //operands that access two variables. @@ -725,6 +732,28 @@ void add_4vector_operand( const char *d1, const char *d2, const char *d3, const if ( s4 == null_op ) s4 = NULL; g_operands.push_back( operand_info(s1,s2,s3,s4) ); } +void add_8vector_operand( const char *d1, const char *d2, const char *d3, const char *d4,const char *d5,const char *d6,const char *d7,const char *d8 ) +{ + PTX_PARSE_DPRINTF("add_8vector_operand"); + const symbol *s1 = g_current_symbol_table->lookup(d1); + const symbol *s2 = g_current_symbol_table->lookup(d2); + const symbol *s3 = g_current_symbol_table->lookup(d3); + const symbol *s4 = g_current_symbol_table->lookup(d4); + const symbol *s5 = g_current_symbol_table->lookup(d5); + const symbol *s6 = g_current_symbol_table->lookup(d6); + const symbol *s7 = g_current_symbol_table->lookup(d7); + const symbol *s8 = g_current_symbol_table->lookup(d8); + parse_assert( s1 != NULL && s2 != NULL && s3 != NULL && s4 != NULL && s5 !=NULL && s6 !=NULL && s7 !=NULL && s8 !=NULL, "v4 component(s) missing declarations."); + const symbol *null_op = g_current_symbol_table->lookup("_"); + if ( s2 == null_op ) s2 = NULL; + if ( s3 == null_op ) s3 = NULL; + if ( s4 == null_op ) s4 = NULL; + if ( s5 == null_op ) s5 = NULL; + if ( s6 == null_op ) s6 = NULL; + if ( s7 == null_op ) s7 = NULL; + if ( s8 == null_op ) s8 = NULL; + g_operands.push_back( operand_info(s1,s2,s3,s4,s5,s6,s7,s8) ); +} void add_builtin_operand( int builtin, int dim_modifier ) { diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index 32f3903..8094b43 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -57,7 +57,9 @@ void add_1vector_operand( const char *d1 ); void add_2vector_operand( const char *d1, const char *d2 ); void add_3vector_operand( const char *d1, const char *d2, const char *d3 ); void add_4vector_operand( const char *d1, const char *d2, const char *d3, const char *d4 ); +void add_8vector_operand( const char *d1, const char *d2, const char *d3, const char *d4 ,const char *d5,const char *d6,const char *d7,const char *d8); void add_option(int option ); +void add_wmma_option(int option ); void add_builtin_operand( int builtin, int dim_modifier ); void add_memory_operand( ); void add_literal_int( int value ); |