|
- //
- // Generated by NVIDIA NVVM Compiler
- //
- // Compiler Build ID: CL-26907403
- // Cuda compilation tools, release 10.1, V10.1.243
- // Based on LLVM 3.4svn
- //
-
- .version 6.4
- .target sm_60
- .address_size 64
-
- // .globl Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0
- // _ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E18input_0_red_shared has been demoted
- // _ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E8red_buf0 has been demoted
-
- .visible .entry Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0(
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_0,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_1,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_2,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_3,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_4,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_5,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_6,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_7,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_8,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_9,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_10,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_11,
- .param .u64 Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_12
- )
- {
- .reg .pred %p<17>;
- .reg .f32 %f<58>;
- .reg .b32 %r<27>;
- .reg .b64 %rd<29>;
- // demoted variable
- .shared .align 4 .b8 _ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E18input_0_red_shared[4];
- // demoted variable
- .shared .align 4 .b8 _ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E8red_buf0[2048];
-
- ld.param.u64 %rd3, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_0];
- ld.param.u64 %rd6, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_1];
- ld.param.u64 %rd7, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_2];
- ld.param.u64 %rd8, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_3];
- ld.param.u64 %rd9, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_4];
- ld.param.u64 %rd10, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_5];
- ld.param.u64 %rd4, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_6];
- ld.param.u64 %rd11, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_7];
- ld.param.u64 %rd12, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_8];
- ld.param.u64 %rd13, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_9];
- ld.param.u64 %rd14, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_10];
- ld.param.u64 %rd15, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_11];
- ld.param.u64 %rd5, [Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0_param_12];
- cvta.to.global.u64 %rd16, %rd7;
- ld.global.nc.f32 %f13, [%rd16];
- mul.f32 %f1, %f13, 0f38555555;
- cvta.to.global.u64 %rd17, %rd12;
- st.global.f32 [%rd17], %f1;
- cvta.to.global.u64 %rd18, %rd9;
- ld.global.nc.f32 %f14, [%rd18];
- mul.f32 %f2, %f14, 0f37D55558;
- cvta.to.global.u64 %rd19, %rd14;
- st.global.f32 [%rd19], %f2;
- cvta.to.global.u64 %rd20, %rd10;
- ld.global.nc.f32 %f15, [%rd20];
- mul.f32 %f3, %f15, 0f37D55558;
- cvta.to.global.u64 %rd21, %rd15;
- st.global.f32 [%rd21], %f3;
- cvta.to.global.u64 %rd22, %rd8;
- ld.global.nc.f32 %f16, [%rd22];
- mul.f32 %f4, %f16, 0f38555555;
- cvta.to.global.u64 %rd23, %rd13;
- st.global.f32 [%rd23], %f4;
- cvta.to.global.u64 %rd24, %rd6;
- ld.global.nc.f32 %f17, [%rd24];
- mul.f32 %f5, %f17, 0f3A91A2B3;
- cvta.to.global.u64 %rd25, %rd11;
- st.global.f32 [%rd25], %f5;
- mov.u32 %r1, %tid.x;
- setp.ne.s32 %p2, %r1, 0;
- @%p2 bra BB0_2;
-
- mov.u32 %r4, 0;
- st.shared.u32 [_ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E18input_0_red_shared], %r4;
-
- BB0_2:
- cvta.to.global.u64 %rd1, %rd3;
- bar.sync 0;
- mul.wide.s32 %rd26, %r1, 4;
- add.s64 %rd2, %rd1, %rd26;
- mov.f32 %f55, 0f00000000;
- setp.gt.s32 %p3, %r1, 899;
- mov.f32 %f56, %f55;
- @%p3 bra BB0_4;
-
- ld.global.nc.f32 %f20, [%rd2];
- add.f32 %f55, %f20, 0f00000000;
- sub.f32 %f56, %f55, %f20;
-
- BB0_4:
- add.s32 %r5, %r1, 512;
- setp.gt.s32 %p4, %r5, 899;
- @%p4 bra BB0_6;
-
- ld.global.nc.f32 %f21, [%rd2+2048];
- sub.f32 %f22, %f21, %f56;
- add.f32 %f55, %f55, %f22;
-
- BB0_6:
- mov.u32 %r6, %tid.y;
- mov.u32 %r7, %ntid.x;
- mad.lo.s32 %r2, %r6, %r7, %r1;
- shl.b32 %r8, %r2, 2;
- mov.u32 %r9, _ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E8red_buf0;
- add.s32 %r3, %r9, %r8;
- st.shared.f32 [%r3], %f55;
- bar.sync 0;
- setp.gt.s32 %p5, %r2, 255;
- @%p5 bra BB0_8;
-
- ld.shared.f32 %f23, [%r3];
- ld.shared.f32 %f24, [%r3+1024];
- add.f32 %f25, %f23, %f24;
- st.shared.f32 [%r3], %f25;
-
- BB0_8:
- bar.sync 0;
- setp.gt.s32 %p6, %r2, 127;
- @%p6 bra BB0_10;
-
- ld.shared.f32 %f26, [%r3];
- ld.shared.f32 %f27, [%r3+512];
- add.f32 %f28, %f26, %f27;
- st.shared.f32 [%r3], %f28;
-
- BB0_10:
- bar.sync 0;
- setp.gt.s32 %p7, %r2, 63;
- @%p7 bra BB0_12;
-
- ld.shared.f32 %f29, [%r3];
- ld.shared.f32 %f30, [%r3+256];
- add.f32 %f31, %f29, %f30;
- st.shared.f32 [%r3], %f31;
-
- BB0_12:
- bar.sync 0;
- setp.gt.s32 %p8, %r2, 31;
- @%p8 bra BB0_14;
-
- ld.shared.f32 %f32, [%r3];
- ld.shared.f32 %f33, [%r3+128];
- add.f32 %f34, %f32, %f33;
- st.shared.f32 [%r3], %f34;
-
- BB0_14:
- setp.lt.s32 %p1, %r2, 32;
- bar.sync 0;
- @!%p1 bra BB0_17;
- bra.uni BB0_15;
-
- BB0_15:
- ld.shared.f32 %f35, [%r3];
- mov.b32 %r10, %f35;
- mov.u32 %r11, 2;
- mov.u32 %r12, 31;
- mov.u32 %r13, 16;
- mov.u32 %r14, -1;
- shfl.sync.down.b32 %r15|%p9, %r10, %r13, %r12, %r14;
- mov.b32 %f36, %r15;
- add.f32 %f37, %f35, %f36;
- mov.b32 %r16, %f37;
- mov.u32 %r17, 8;
- shfl.sync.down.b32 %r18|%p10, %r16, %r17, %r12, %r14;
- mov.b32 %f38, %r18;
- add.f32 %f39, %f37, %f38;
- mov.b32 %r19, %f39;
- mov.u32 %r20, 4;
- shfl.sync.down.b32 %r21|%p11, %r19, %r20, %r12, %r14;
- mov.b32 %f40, %r21;
- add.f32 %f41, %f39, %f40;
- mov.b32 %r22, %f41;
- shfl.sync.down.b32 %r23|%p12, %r22, %r11, %r12, %r14;
- mov.b32 %f42, %r23;
- add.f32 %f43, %f41, %f42;
- mov.b32 %r24, %f43;
- mov.u32 %r25, 1;
- shfl.sync.down.b32 %r26|%p13, %r24, %r25, %r12, %r14;
- mov.b32 %f44, %r26;
- add.f32 %f12, %f43, %f44;
- setp.ne.s32 %p14, %r2, 0;
- @%p14 bra BB0_17;
-
- st.shared.f32 [_ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E8red_buf0], %f12;
-
- BB0_17:
- bar.sync 0;
- setp.ne.s32 %p15, %r2, 0;
- @%p15 bra BB0_19;
-
- ld.shared.f32 %f45, [_ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E18input_0_red_shared];
- ld.shared.f32 %f46, [_ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E8red_buf0];
- add.f32 %f47, %f45, %f46;
- st.shared.f32 [_ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E18input_0_red_shared], %f47;
-
- BB0_19:
- bar.sync 0;
- setp.ne.s32 %p16, %r1, 387;
- @%p16 bra BB0_21;
-
- ld.shared.f32 %f48, [_ZZ93Fused_ReduceSum_Mul_Mul_Mul_Mul_Mul_Mul_Add_Add_Add_Add_Add_split_1940226729548372996_kernel0E18input_0_red_shared];
- mul.f32 %f49, %f48, 0f3A91A2B3;
- cvta.to.global.u64 %rd27, %rd4;
- st.global.f32 [%rd27], %f49;
- add.f32 %f50, %f49, %f5;
- add.f32 %f51, %f50, %f1;
- add.f32 %f52, %f51, %f4;
- add.f32 %f53, %f52, %f2;
- add.f32 %f54, %f53, %f3;
- cvta.to.global.u64 %rd28, %rd5;
- st.global.f32 [%rd28], %f54;
-
- BB0_21:
- bar.sync 0;
- bar.sync 0;
- ret;
- }
-
|