|
- //
- // 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_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0
-
- .visible .entry Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0(
- .param .u64 Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_0,
- .param .u64 Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_1,
- .param .u64 Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_2,
- .param .u64 Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_3
- )
- {
- .reg .pred %p<13>;
- .reg .f32 %f<91>;
- .reg .b32 %r<13>;
- .reg .b64 %rd<19>;
-
-
- ld.param.u64 %rd5, [Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_0];
- ld.param.u64 %rd6, [Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_1];
- ld.param.u64 %rd7, [Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_2];
- ld.param.u64 %rd8, [Fused_MaximumGrad_Mul_Add_fusion_16806538634415810313_kernel0_param_3];
- cvta.to.global.u64 %rd1, %rd8;
- cvta.to.global.u64 %rd2, %rd5;
- cvta.to.global.u64 %rd3, %rd6;
- cvta.to.global.u64 %rd4, %rd7;
- mov.u32 %r1, %tid.x;
- setp.gt.s32 %p1, %r1, 1015;
- @%p1 bra BB0_3;
-
- mov.u32 %r3, %ctaid.x;
- shl.b32 %r2, %r1, 2;
- mad.lo.s32 %r4, %r3, 4064, %r2;
- mul.wide.s32 %rd9, %r4, 4;
- add.s64 %rd10, %rd4, %rd9;
- ld.global.nc.v4.f32 {%f1, %f2, %f3, %f4}, [%rd10];
- add.s64 %rd11, %rd3, %rd9;
- ld.global.nc.v4.f32 {%f9, %f10, %f11, %f12}, [%rd11];
- add.s64 %rd12, %rd2, %rd9;
- ld.global.nc.v4.f32 {%f17, %f18, %f19, %f20}, [%rd12];
- setp.le.f32 %p2, %f9, %f17;
- mul.f32 %f25, %f1, 0f3E4CCCCD;
- selp.f32 %f26, %f25, 0f00000000, %p2;
- selp.u32 %r5, 1, 0, %p2;
- cvt.rn.f32.s32 %f27, %r5;
- mov.f32 %f28, 0f3F800000;
- sub.f32 %f29, %f28, %f27;
- setp.le.f32 %p3, %f10, %f18;
- mul.f32 %f30, %f2, 0f3E4CCCCD;
- selp.f32 %f31, %f30, 0f00000000, %p3;
- selp.u32 %r6, 1, 0, %p3;
- cvt.rn.f32.s32 %f32, %r6;
- sub.f32 %f33, %f28, %f32;
- setp.le.f32 %p4, %f11, %f19;
- mul.f32 %f34, %f3, 0f3E4CCCCD;
- selp.f32 %f35, %f34, 0f00000000, %p4;
- selp.u32 %r7, 1, 0, %p4;
- cvt.rn.f32.s32 %f36, %r7;
- sub.f32 %f37, %f28, %f36;
- setp.le.f32 %p5, %f12, %f20;
- mul.f32 %f38, %f4, 0f3E4CCCCD;
- selp.f32 %f39, %f38, 0f00000000, %p5;
- selp.u32 %r8, 1, 0, %p5;
- cvt.rn.f32.s32 %f40, %r8;
- sub.f32 %f41, %f28, %f40;
- add.s64 %rd13, %rd1, %rd9;
- fma.rn.f32 %f42, %f4, %f41, %f39;
- fma.rn.f32 %f43, %f3, %f37, %f35;
- fma.rn.f32 %f44, %f2, %f33, %f31;
- fma.rn.f32 %f45, %f1, %f29, %f26;
- st.global.v4.f32 [%rd13], {%f45, %f44, %f43, %f42};
- setp.ne.s32 %p6, %r3, 0;
- setp.gt.s32 %p7, %r1, 71;
- or.pred %p8, %p6, %p7;
- @%p8 bra BB0_3;
-
- mul.wide.s32 %rd14, %r2, 4;
- add.s64 %rd15, %rd4, %rd14;
- ld.global.nc.v4.f32 {%f46, %f47, %f48, %f49}, [%rd15+1966976];
- add.s64 %rd16, %rd3, %rd14;
- ld.global.nc.v4.f32 {%f54, %f55, %f56, %f57}, [%rd16+1966976];
- add.s64 %rd17, %rd2, %rd14;
- ld.global.nc.v4.f32 {%f62, %f63, %f64, %f65}, [%rd17+1966976];
- setp.le.f32 %p9, %f54, %f62;
- mul.f32 %f70, %f46, 0f3E4CCCCD;
- selp.f32 %f71, %f70, 0f00000000, %p9;
- selp.u32 %r9, 1, 0, %p9;
- cvt.rn.f32.s32 %f72, %r9;
- sub.f32 %f74, %f28, %f72;
- setp.le.f32 %p10, %f55, %f63;
- mul.f32 %f75, %f47, 0f3E4CCCCD;
- selp.f32 %f76, %f75, 0f00000000, %p10;
- selp.u32 %r10, 1, 0, %p10;
- cvt.rn.f32.s32 %f77, %r10;
- sub.f32 %f78, %f28, %f77;
- setp.le.f32 %p11, %f56, %f64;
- mul.f32 %f79, %f48, 0f3E4CCCCD;
- selp.f32 %f80, %f79, 0f00000000, %p11;
- selp.u32 %r11, 1, 0, %p11;
- cvt.rn.f32.s32 %f81, %r11;
- sub.f32 %f82, %f28, %f81;
- setp.le.f32 %p12, %f57, %f65;
- mul.f32 %f83, %f49, 0f3E4CCCCD;
- selp.f32 %f84, %f83, 0f00000000, %p12;
- selp.u32 %r12, 1, 0, %p12;
- cvt.rn.f32.s32 %f85, %r12;
- sub.f32 %f86, %f28, %f85;
- add.s64 %rd18, %rd1, %rd14;
- fma.rn.f32 %f87, %f49, %f86, %f84;
- fma.rn.f32 %f88, %f48, %f82, %f80;
- fma.rn.f32 %f89, %f47, %f78, %f76;
- fma.rn.f32 %f90, %f46, %f74, %f71;
- st.global.v4.f32 [%rd18+1966976], {%f90, %f89, %f88, %f87};
-
- BB0_3:
- ret;
- }
-
|