|
- //
- // 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_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0
-
- .visible .entry Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0(
- .param .u64 Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_0,
- .param .u64 Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_1,
- .param .u64 Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_2,
- .param .u64 Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_3,
- .param .u64 Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_4
- )
- {
- .reg .pred %p<5>;
- .reg .f32 %f<58>;
- .reg .b32 %r<10>;
- .reg .b64 %rd<17>;
-
-
- ld.param.u64 %rd1, [Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_0];
- ld.param.u64 %rd2, [Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_1];
- ld.param.u64 %rd3, [Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_2];
- ld.param.u64 %rd4, [Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_3];
- ld.param.u64 %rd5, [Fused_Add_MaximumGrad_Mul_Add_fusion_3028261687580406729_kernel0_param_4];
- cvta.to.global.u64 %rd6, %rd3;
- cvta.to.global.u64 %rd7, %rd4;
- cvta.to.global.u64 %rd8, %rd1;
- cvta.to.global.u64 %rd9, %rd2;
- mov.u32 %r1, %ctaid.x;
- shl.b32 %r2, %r1, 12;
- mov.u32 %r3, %tid.x;
- shl.b32 %r4, %r3, 2;
- add.s32 %r5, %r4, %r2;
- mul.wide.s32 %rd10, %r5, 4;
- add.s64 %rd11, %rd9, %rd10;
- ld.global.nc.v4.f32 {%f1, %f2, %f3, %f4}, [%rd11];
- add.s64 %rd12, %rd8, %rd10;
- ld.global.nc.v4.f32 {%f9, %f10, %f11, %f12}, [%rd12];
- add.s64 %rd13, %rd7, %rd10;
- ld.global.nc.v4.f32 {%f17, %f18, %f19, %f20}, [%rd13];
- add.s64 %rd14, %rd6, %rd10;
- ld.global.nc.v4.f32 {%f25, %f26, %f27, %f28}, [%rd14];
- setp.le.f32 %p1, %f17, %f25;
- add.f32 %f33, %f9, %f1;
- mul.f32 %f34, %f33, 0f3E4CCCCD;
- selp.f32 %f35, %f34, 0f00000000, %p1;
- selp.u32 %r6, 1, 0, %p1;
- cvt.rn.f32.s32 %f36, %r6;
- mov.f32 %f37, 0f3F800000;
- sub.f32 %f38, %f37, %f36;
- setp.le.f32 %p2, %f18, %f26;
- add.f32 %f39, %f10, %f2;
- mul.f32 %f40, %f39, 0f3E4CCCCD;
- selp.f32 %f41, %f40, 0f00000000, %p2;
- selp.u32 %r7, 1, 0, %p2;
- cvt.rn.f32.s32 %f42, %r7;
- sub.f32 %f43, %f37, %f42;
- setp.le.f32 %p3, %f19, %f27;
- add.f32 %f44, %f11, %f3;
- mul.f32 %f45, %f44, 0f3E4CCCCD;
- selp.f32 %f46, %f45, 0f00000000, %p3;
- selp.u32 %r8, 1, 0, %p3;
- cvt.rn.f32.s32 %f47, %r8;
- sub.f32 %f48, %f37, %f47;
- setp.le.f32 %p4, %f20, %f28;
- add.f32 %f49, %f12, %f4;
- mul.f32 %f50, %f49, 0f3E4CCCCD;
- selp.f32 %f51, %f50, 0f00000000, %p4;
- selp.u32 %r9, 1, 0, %p4;
- cvt.rn.f32.s32 %f52, %r9;
- sub.f32 %f53, %f37, %f52;
- cvta.to.global.u64 %rd15, %rd5;
- add.s64 %rd16, %rd15, %rd10;
- fma.rn.f32 %f54, %f49, %f53, %f51;
- fma.rn.f32 %f55, %f44, %f48, %f46;
- fma.rn.f32 %f56, %f39, %f43, %f41;
- fma.rn.f32 %f57, %f33, %f38, %f35;
- st.global.v4.f32 [%rd16], {%f57, %f56, %f55, %f54};
- ret;
- }
-
|