diff options
Diffstat (limited to 'cuda-kernels/tensor_core_ptx')
| -rw-r--r-- | cuda-kernels/tensor_core_ptx | 171 |
1 files changed, 0 insertions, 171 deletions
diff --git a/cuda-kernels/tensor_core_ptx b/cuda-kernels/tensor_core_ptx deleted file mode 100644 index 36074cb..0000000 --- a/cuda-kernels/tensor_core_ptx +++ /dev/null @@ -1,171 +0,0 @@ -// -// Generated by NVIDIA NVVM Compiler -// -// Compiler Build ID: CL-22781540 -// Cuda compilation tools, release 9.0, V9.0.176 -// Based on LLVM 3.4svn -// - -.version 6.0 -.target sm_70 -.address_size 64 - - // .globl _Z12wmma_exampleP6__halfS0_Pfiiiff -.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]; - // inline asm - mov.u32 %r6, %clock; - // inline asm - 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; - // inline asm - mov.u32 %r34, %clock; - // inline asm - sub.s32 %r36, %r34, %r6; - st.local.u32 [%rd9], %r36; - mov.u64 %rd15, $str; - cvta.global.u64 %rd16, %rd15; - // Callseq Start 0 - { - .reg .b32 temp_param_reg; - // <end>} - .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]; - - //{ - }// Callseq End 0 - ret; -} - - // .globl _Z17convertFp32ToFp16P6__halfPfi -.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]; - // inline asm - { cvt.rn.f16.f32 %rs1, %f1;} - - // inline asm - 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; -} - - |