1 files changed, 0 insertions, 170 deletions
diff --git a/cuda-kernels/_cuobjdump_1.ptx b/cuda-kernels/_cuobjdump_1.ptx
deleted file mode 100644
index 3453f4a..0000000
--- a/cuda-kernels/_cuobjdump_1.ptx
+++ /dev/null
@@ -1,170 +0,0 @@
-
-
-
-
-
-
-
-.version 6.0
-.target sm_70
-.address_size 64
-
-
-.extern .func (.param .b32 func_retval0) vprintf
-(
-.param .b64 vprintf_param_0,
-.param .b64 vprintf_param_1
-)
-;
-.global .align 16 .b8 $str[9] = {99, 108, 111, 99, 107, 61, 37, 100, 0};
-
-.visible .entry _Z12wmma_exampleP6__halfS0_Pfiiiff(
-.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_0,
-.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_1,
-.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_2,
-.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_3,
-.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_4,
-.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_5,
-.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_6,
-.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_7
-)
-{
-.local .align 8 .b8 __local_depot0[8];
-.reg .b64 %SP;
-.reg .b64 %SPL;
-.reg .pred %p<6>;
-.reg .f32 %f<34>;
-.reg .b32 %r<38>;
-.reg .b64 %rd<18>;
-
-
-mov.u64 %rd17, __local_depot0;
-cvta.local.u64 %SP, %rd17;
-ld.param.u64 %rd1, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0];
-ld.param.u64 %rd2, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1];
-ld.param.u64 %rd3, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2];
-ld.param.u32 %r4, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3];
-ld.param.u32 %r7, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4];
-ld.param.u32 %r5, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5];
-
-	mov.u32 %r6, %clock;
-
-	mov.u32 %r8, %ntid.x;
-mov.u32 %r9, %ctaid.x;
-mov.u32 %r10, %tid.x;
-mad.lo.s32 %r11, %r8, %r9, %r10;
-mov.u32 %r12, WARP_SZ;
-div.u32 %r13, %r11, %r12;
-mov.u32 %r14, %ntid.y;
-mov.u32 %r15, %ctaid.y;
-mov.u32 %r16, %tid.y;
-mad.lo.s32 %r17, %r14, %r15, %r16;
-shl.b32 %r2, %r13, 4;
-shl.b32 %r3, %r17, 4;
-setp.lt.s32	%p1, %r2, %r4;
-setp.gt.s32	%p2, %r5, 0;
-and.pred %p3, %p1, %p2;
-setp.lt.s32	%p4, %r3, %r7;
-and.pred %p5, %p3, %p4;
-mov.f32 %f26, 0f00000000;
-mov.f32 %f27, %f26;
-mov.f32 %f28, %f26;
-mov.f32 %f29, %f26;
-mov.f32 %f30, %f26;
-mov.f32 %f31, %f26;
-mov.f32 %f32, %f26;
-mov.f32 %f33, %f26;
-@!%p5 bra BB0_2;
-bra.uni BB0_1;
-
-BB0_1:
-mul.wide.s32 %rd4, %r2, 2;
-add.s64 %rd5, %rd1, %rd4;
-wmma.load.a.sync.row.m16n16k16.f16 {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, [%rd5], %r4;
-mul.wide.s32 %rd6, %r3, 2;
-add.s64 %rd7, %rd2, %rd6;
-wmma.load.b.sync.col.m16n16k16.f16 {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, [%rd7], %r5;
-mov.f32 %f25, 0f00000000;
-wmma.mma.sync.row.col.m16n16k16.f32.f32 {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, {%f25, %f25, %f25, %f25, %f25, %f25, %f25, %f25};
-
-BB0_2:
-add.u64 %rd8, %SP, 0;
-cvta.to.local.u64 %rd9, %rd8;
-mul.lo.s32 %r35, %r3, %r4;
-cvt.s64.s32	%rd10, %r35;
-cvt.s64.s32	%rd11, %r2;
-add.s64 %rd12, %rd10, %rd11;
-shl.b64 %rd13, %rd12, 2;
-add.s64 %rd14, %rd3, %rd13;
-wmma.store.d.sync.col.m16n16k16.f32 [%rd14], {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, %r4;
-
-	mov.u32 %r34, %clock;
-
-	sub.s32 %r36, %r34, %r6;
-st.local.u32 [%rd9], %r36;
-mov.u64 %rd15, $str;
-cvta.global.u64 %rd16, %rd15;
-
-	{
-.reg .b32 temp_param_reg;
-
-	.param .b64 param0;
-st.param.b64	[param0+0], %rd16;
-.param .b64 param1;
-st.param.b64	[param1+0], %rd8;
-.param .b32 retval0;
-call.uni (retval0), 
-vprintf, 
-(
-param0, 
-param1
-);
-ld.param.b32	%r37, [retval0+0];
-
-
-	}
-	ret;
-}
-
-
-.visible .entry _Z17convertFp32ToFp16P6__halfPfi(
-.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_0,
-.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_1,
-.param .u32 _Z17convertFp32ToFp16P6__halfPfi_param_2
-)
-{
-.reg .pred %p<2>;
-.reg .b16 %rs<2>;
-.reg .f32 %f<2>;
-.reg .b32 %r<6>;
-.reg .b64 %rd<9>;
-
-
-ld.param.u64 %rd1, [_Z17convertFp32ToFp16P6__halfPfi_param_0];
-ld.param.u64 %rd2, [_Z17convertFp32ToFp16P6__halfPfi_param_1];
-ld.param.u32 %r2, [_Z17convertFp32ToFp16P6__halfPfi_param_2];
-mov.u32 %r3, %ntid.x;
-mov.u32 %r4, %ctaid.x;
-mov.u32 %r5, %tid.x;
-mad.lo.s32 %r1, %r4, %r3, %r5;
-setp.ge.s32	%p1, %r1, %r2;
-@%p1 bra BB1_2;
-
-cvta.to.global.u64 %rd3, %rd2;
-mul.wide.s32 %rd4, %r1, 4;
-add.s64 %rd5, %rd3, %rd4;
-ld.global.f32 %f1, [%rd5];
-
-	{ cvt.rn.f16.f32 %rs1, %f1;}
-
-
-	cvta.to.global.u64 %rd6, %rd1;
-mul.wide.s32 %rd7, %r1, 2;
-add.s64 %rd8, %rd6, %rd7;
-st.global.u16 [%rd8], %rs1;
-
-BB1_2:
-ret;
-}
-
-