|
- //
- // 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_Mul_Maximum_fusion_6621062658143142625_kernel0
-
- .visible .entry Fused_Mul_Maximum_fusion_6621062658143142625_kernel0(
- .param .u64 Fused_Mul_Maximum_fusion_6621062658143142625_kernel0_param_0,
- .param .u64 Fused_Mul_Maximum_fusion_6621062658143142625_kernel0_param_1
- )
- {
- .reg .f32 %f<65>;
- .reg .b32 %r<6>;
- .reg .b64 %rd<8>;
-
-
- ld.param.u64 %rd1, [Fused_Mul_Maximum_fusion_6621062658143142625_kernel0_param_0];
- ld.param.u64 %rd2, [Fused_Mul_Maximum_fusion_6621062658143142625_kernel0_param_1];
- cvta.to.global.u64 %rd3, %rd2;
- cvta.to.global.u64 %rd4, %rd1;
- mov.u32 %r1, %ctaid.x;
- shl.b32 %r2, %r1, 13;
- mov.u32 %r3, %tid.x;
- shl.b32 %r4, %r3, 2;
- add.s32 %r5, %r2, %r4;
- mul.wide.s32 %rd5, %r5, 4;
- add.s64 %rd6, %rd4, %rd5;
- ld.global.nc.v4.f32 {%f1, %f2, %f3, %f4}, [%rd6];
- mul.f32 %f9, %f1, 0f3E4CCCCD;
- mul.f32 %f10, %f2, 0f3E4CCCCD;
- mul.f32 %f11, %f3, 0f3E4CCCCD;
- mul.f32 %f12, %f4, 0f3E4CCCCD;
- add.s64 %rd7, %rd3, %rd5;
- max.f32 %f13, %f12, %f4;
- max.f32 %f14, %f11, %f3;
- max.f32 %f15, %f10, %f2;
- max.f32 %f16, %f9, %f1;
- st.global.v4.f32 [%rd7], {%f16, %f15, %f14, %f13};
- ld.global.nc.v4.f32 {%f17, %f18, %f19, %f20}, [%rd6+8192];
- mul.f32 %f25, %f17, 0f3E4CCCCD;
- mul.f32 %f26, %f18, 0f3E4CCCCD;
- mul.f32 %f27, %f19, 0f3E4CCCCD;
- mul.f32 %f28, %f20, 0f3E4CCCCD;
- max.f32 %f29, %f28, %f20;
- max.f32 %f30, %f27, %f19;
- max.f32 %f31, %f26, %f18;
- max.f32 %f32, %f25, %f17;
- st.global.v4.f32 [%rd7+8192], {%f32, %f31, %f30, %f29};
- ld.global.nc.v4.f32 {%f33, %f34, %f35, %f36}, [%rd6+16384];
- mul.f32 %f41, %f33, 0f3E4CCCCD;
- mul.f32 %f42, %f34, 0f3E4CCCCD;
- mul.f32 %f43, %f35, 0f3E4CCCCD;
- mul.f32 %f44, %f36, 0f3E4CCCCD;
- max.f32 %f45, %f44, %f36;
- max.f32 %f46, %f43, %f35;
- max.f32 %f47, %f42, %f34;
- max.f32 %f48, %f41, %f33;
- st.global.v4.f32 [%rd7+16384], {%f48, %f47, %f46, %f45};
- ld.global.nc.v4.f32 {%f49, %f50, %f51, %f52}, [%rd6+24576];
- mul.f32 %f57, %f49, 0f3E4CCCCD;
- mul.f32 %f58, %f50, 0f3E4CCCCD;
- mul.f32 %f59, %f51, 0f3E4CCCCD;
- mul.f32 %f60, %f52, 0f3E4CCCCD;
- max.f32 %f61, %f60, %f52;
- max.f32 %f62, %f59, %f51;
- max.f32 %f63, %f58, %f50;
- max.f32 %f64, %f57, %f49;
- st.global.v4.f32 [%rd7+24576], {%f64, %f63, %f62, %f61};
- ret;
- }
-
|