another tentative

Author: N

examples/multi_tensor.ptx

@@ -0,0 +1,129 @@
+.version 8.0
+.target sm_52
+.address_size 64
+
+
+
+.visible .entry _Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i(
+.param .u64 _Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_0,
+.param .u64 _Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_1,
+.param .u64 _Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_2,
+.param .u32 _Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_3
+)
+{
+.reg .pred %p<4>;
+.reg .f32 %f<4>;
+.reg .b32 %r<15>;
+.reg .b64 %rd<21>;
+
+
+ld.param.u64 %rd7, [_Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_0];
+ld.param.u64 %rd8, [_Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_1];
+ld.param.u64 %rd9, [_Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_2];
+ld.param.u32 %r7, [_Z36multi_tensor_add_kernel_nested_arrayPPfS0_S0_i_param_3];
+cvta.to.global.u64 %rd1, %rd9;
+cvta.to.global.u64 %rd2, %rd8;
+cvta.to.global.u64 %rd3, %rd7;
+mov.u32 %r9, %ntid.x;
+mov.u32 %r10, %ctaid.x;
+mov.u32 %r11, %tid.x;
+mad.lo.s32 %r1, %r10, %r9, %r11;
+mov.u32 %r12, %nctaid.x;
+mul.lo.s32 %r2, %r9, %r12;
+mov.u32 %r13, 0;
+
+$L__BB0_1:
+setp.ge.s32 %p1, %r1, %r7;
+@%p1 bra $L__BB0_4;
+
+mul.wide.s32 %rd10, %r13, 8;
+add.s64 %rd4, %rd3, %rd10;
+add.s64 %rd5, %rd2, %rd10;
+add.s64 %rd6, %rd1, %rd10;
+mov.u32 %r14, %r1;
+
+$L__BB0_3:
+ld.global.u64 %rd11, [%rd4];
+cvta.to.global.u64 %rd12, %rd11;
+mul.wide.s32 %rd13, %r14, 4;
+add.s64 %rd14, %rd12, %rd13;
+ld.global.u64 %rd15, [%rd5];
+cvta.to.global.u64 %rd16, %rd15;
+add.s64 %rd17, %rd16, %rd13;
+ld.global.f32 %f1, [%rd17];
+ld.global.f32 %f2, [%rd14];
+add.ftz.f32 %f3, %f2, %f1;
+ld.global.u64 %rd18, [%rd6];
+cvta.to.global.u64 %rd19, %rd18;
+add.s64 %rd20, %rd19, %rd13;
+st.global.f32 [%rd20], %f3;
+add.s32 %r14, %r14, %r2;
+setp.lt.s32 %p2, %r14, %r7;
+@%p2 bra $L__BB0_3;
+
+$L__BB0_4:
+add.s32 %r13, %r13, 1;
+setp.lt.u32 %p3, %r13, 1024;
+@%p3 bra $L__BB0_1;
+
+ret;
+
+}
+
+.visible .entry _Z33multi_tensor_add_kernel_meta_dataP18TensorListMetaDatai(
+.param .u64 _Z33multi_tensor_add_kernel_meta_dataP18TensorListMetaDatai_param_0,
+.param .u32 _Z33multi_tensor_add_kernel_meta_dataP18TensorListMetaDatai_param_1
+)
+{
+.reg .pred %p<4>;
+.reg .f32 %f<4>;
+.reg .b32 %r<15>;
+.reg .b64 %rd<15>;
+
+
+ld.param.u64 %rd5, [_Z33multi_tensor_add_kernel_meta_dataP18TensorListMetaDatai_param_0];
+ld.param.u32 %r7, [_Z33multi_tensor_add_kernel_meta_dataP18TensorListMetaDatai_param_1];
+cvta.to.global.u64 %rd1, %rd5;
+mov.u32 %r9, %ntid.x;
+mov.u32 %r10, %ctaid.x;
+mov.u32 %r11, %tid.x;
+mad.lo.s32 %r1, %r10, %r9, %r11;
+mov.u32 %r12, %nctaid.x;
+mul.lo.s32 %r2, %r9, %r12;
+mov.u32 %r13, 0;
+
+$L__BB1_1:
+setp.ge.s32 %p1, %r1, %r7;
+@%p1 bra $L__BB1_4;
+
+mul.wide.s32 %rd6, %r13, 8;
+add.s64 %rd7, %rd1, %rd6;
+ld.global.u64 %rd8, [%rd7];
+ld.global.u64 %rd9, [%rd7+8192];
+ld.global.u64 %rd10, [%rd7+16384];
+cvta.to.global.u64 %rd2, %rd10;
+cvta.to.global.u64 %rd3, %rd9;
+cvta.to.global.u64 %rd4, %rd8;
+mov.u32 %r14, %r1;
+
+$L__BB1_3:
+mul.wide.s32 %rd11, %r14, 4;
+add.s64 %rd12, %rd4, %rd11;
+add.s64 %rd13, %rd3, %rd11;
+ld.global.f32 %f1, [%rd13];
+ld.global.f32 %f2, [%rd12];
+add.ftz.f32 %f3, %f2, %f1;
+add.s64 %rd14, %rd2, %rd11;
+st.global.f32 [%rd14], %f3;
+add.s32 %r14, %r14, %r2;
+setp.lt.s32 %p2, %r14, %r7;
+@%p2 bra $L__BB1_3;
+
+$L__BB1_4:
+add.s32 %r13, %r13, 1;
+setp.lt.u32 %p3, %r13, 1024;
+@%p3 bra $L__BB1_1;
+
+ret;
+
+}