1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
3 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
5 target triple = "nvptx64-nvidia-cuda"
6 target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
8 define float @fcopysign_f_f(float %a, float %b) {
9 ; CHECK-LABEL: fcopysign_f_f(
11 ; CHECK-NEXT: .reg .f32 %f<4>;
13 ; CHECK-NEXT: // %bb.0:
14 ; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_f_param_0];
15 ; CHECK-NEXT: ld.param.f32 %f2, [fcopysign_f_f_param_1];
16 ; CHECK-NEXT: copysign.f32 %f3, %f2, %f1;
17 ; CHECK-NEXT: st.param.f32 [func_retval0], %f3;
19 %val = call float @llvm.copysign.f32(float %a, float %b)
23 define double @fcopysign_d_d(double %a, double %b) {
24 ; CHECK-LABEL: fcopysign_d_d(
26 ; CHECK-NEXT: .reg .f64 %fd<4>;
28 ; CHECK-NEXT: // %bb.0:
29 ; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_d_param_0];
30 ; CHECK-NEXT: ld.param.f64 %fd2, [fcopysign_d_d_param_1];
31 ; CHECK-NEXT: copysign.f64 %fd3, %fd2, %fd1;
32 ; CHECK-NEXT: st.param.f64 [func_retval0], %fd3;
34 %val = call double @llvm.copysign.f64(double %a, double %b)
38 define float @fcopysign_f_d(float %a, double %b) {
39 ; CHECK-LABEL: fcopysign_f_d(
41 ; CHECK-NEXT: .reg .pred %p<2>;
42 ; CHECK-NEXT: .reg .f32 %f<5>;
43 ; CHECK-NEXT: .reg .b64 %rd<4>;
45 ; CHECK-NEXT: // %bb.0:
46 ; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_d_param_0];
47 ; CHECK-NEXT: abs.f32 %f2, %f1;
48 ; CHECK-NEXT: neg.f32 %f3, %f2;
49 ; CHECK-NEXT: ld.param.u64 %rd1, [fcopysign_f_d_param_1];
50 ; CHECK-NEXT: shr.u64 %rd2, %rd1, 63;
51 ; CHECK-NEXT: and.b64 %rd3, %rd2, 1;
52 ; CHECK-NEXT: setp.eq.b64 %p1, %rd3, 1;
53 ; CHECK-NEXT: selp.f32 %f4, %f3, %f2, %p1;
54 ; CHECK-NEXT: st.param.f32 [func_retval0], %f4;
56 %c = fptrunc double %b to float
57 %val = call float @llvm.copysign.f32(float %a, float %c)
61 define float @fcopysign_f_h(float %a, half %b) {
62 ; CHECK-LABEL: fcopysign_f_h(
64 ; CHECK-NEXT: .reg .pred %p<2>;
65 ; CHECK-NEXT: .reg .b16 %rs<4>;
66 ; CHECK-NEXT: .reg .f32 %f<5>;
68 ; CHECK-NEXT: // %bb.0:
69 ; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_h_param_0];
70 ; CHECK-NEXT: abs.f32 %f2, %f1;
71 ; CHECK-NEXT: neg.f32 %f3, %f2;
72 ; CHECK-NEXT: ld.param.u16 %rs1, [fcopysign_f_h_param_1];
73 ; CHECK-NEXT: shr.u16 %rs2, %rs1, 15;
74 ; CHECK-NEXT: and.b16 %rs3, %rs2, 1;
75 ; CHECK-NEXT: setp.eq.b16 %p1, %rs3, 1;
76 ; CHECK-NEXT: selp.f32 %f4, %f3, %f2, %p1;
77 ; CHECK-NEXT: st.param.f32 [func_retval0], %f4;
79 %c = fpext half %b to float
80 %val = call float @llvm.copysign.f32(float %a, float %c)
84 define double @fcopysign_d_f(double %a, float %b) {
85 ; CHECK-LABEL: fcopysign_d_f(
87 ; CHECK-NEXT: .reg .pred %p<2>;
88 ; CHECK-NEXT: .reg .b32 %r<4>;
89 ; CHECK-NEXT: .reg .f64 %fd<5>;
91 ; CHECK-NEXT: // %bb.0:
92 ; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_f_param_0];
93 ; CHECK-NEXT: abs.f64 %fd2, %fd1;
94 ; CHECK-NEXT: neg.f64 %fd3, %fd2;
95 ; CHECK-NEXT: ld.param.u32 %r1, [fcopysign_d_f_param_1];
96 ; CHECK-NEXT: shr.u32 %r2, %r1, 31;
97 ; CHECK-NEXT: and.b32 %r3, %r2, 1;
98 ; CHECK-NEXT: setp.eq.b32 %p1, %r3, 1;
99 ; CHECK-NEXT: selp.f64 %fd4, %fd3, %fd2, %p1;
100 ; CHECK-NEXT: st.param.f64 [func_retval0], %fd4;
102 %c = fpext float %b to double
103 %val = call double @llvm.copysign.f64(double %a, double %c)
107 define double @fcopysign_d_h(double %a, half %b) {
108 ; CHECK-LABEL: fcopysign_d_h(
110 ; CHECK-NEXT: .reg .pred %p<2>;
111 ; CHECK-NEXT: .reg .b16 %rs<4>;
112 ; CHECK-NEXT: .reg .f64 %fd<5>;
114 ; CHECK-NEXT: // %bb.0:
115 ; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_h_param_0];
116 ; CHECK-NEXT: abs.f64 %fd2, %fd1;
117 ; CHECK-NEXT: neg.f64 %fd3, %fd2;
118 ; CHECK-NEXT: ld.param.u16 %rs1, [fcopysign_d_h_param_1];
119 ; CHECK-NEXT: shr.u16 %rs2, %rs1, 15;
120 ; CHECK-NEXT: and.b16 %rs3, %rs2, 1;
121 ; CHECK-NEXT: setp.eq.b16 %p1, %rs3, 1;
122 ; CHECK-NEXT: selp.f64 %fd4, %fd3, %fd2, %p1;
123 ; CHECK-NEXT: st.param.f64 [func_retval0], %fd4;
125 %c = fpext half %b to double
126 %val = call double @llvm.copysign.f64(double %a, double %c)
131 declare float @llvm.copysign.f32(float, float)
132 declare double @llvm.copysign.f64(double, double)