1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s 3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} 4 5target triple = "nvptx64-nvidia-cuda" 6target 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" 7 8define float @fcopysign_f_f(float %a, float %b) { 9; CHECK-LABEL: fcopysign_f_f( 10; CHECK: { 11; CHECK-NEXT: .reg .f32 %f<4>; 12; CHECK-EMPTY: 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; 18; CHECK-NEXT: ret; 19 %val = call float @llvm.copysign.f32(float %a, float %b) 20 ret float %val 21} 22 23define double @fcopysign_d_d(double %a, double %b) { 24; CHECK-LABEL: fcopysign_d_d( 25; CHECK: { 26; CHECK-NEXT: .reg .f64 %fd<4>; 27; CHECK-EMPTY: 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; 33; CHECK-NEXT: ret; 34 %val = call double @llvm.copysign.f64(double %a, double %b) 35 ret double %val 36} 37 38define float @fcopysign_f_d(float %a, double %b) { 39; CHECK-LABEL: fcopysign_f_d( 40; CHECK: { 41; CHECK-NEXT: .reg .pred %p<2>; 42; CHECK-NEXT: .reg .f32 %f<5>; 43; CHECK-NEXT: .reg .b64 %rd<4>; 44; CHECK-EMPTY: 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; 55; CHECK-NEXT: ret; 56 %c = fptrunc double %b to float 57 %val = call float @llvm.copysign.f32(float %a, float %c) 58 ret float %val 59} 60 61define float @fcopysign_f_h(float %a, half %b) { 62; CHECK-LABEL: fcopysign_f_h( 63; CHECK: { 64; CHECK-NEXT: .reg .pred %p<2>; 65; CHECK-NEXT: .reg .b16 %rs<4>; 66; CHECK-NEXT: .reg .f32 %f<5>; 67; CHECK-EMPTY: 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; 78; CHECK-NEXT: ret; 79 %c = fpext half %b to float 80 %val = call float @llvm.copysign.f32(float %a, float %c) 81 ret float %val 82} 83 84define double @fcopysign_d_f(double %a, float %b) { 85; CHECK-LABEL: fcopysign_d_f( 86; CHECK: { 87; CHECK-NEXT: .reg .pred %p<2>; 88; CHECK-NEXT: .reg .b32 %r<4>; 89; CHECK-NEXT: .reg .f64 %fd<5>; 90; CHECK-EMPTY: 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; 101; CHECK-NEXT: ret; 102 %c = fpext float %b to double 103 %val = call double @llvm.copysign.f64(double %a, double %c) 104 ret double %val 105} 106 107define double @fcopysign_d_h(double %a, half %b) { 108; CHECK-LABEL: fcopysign_d_h( 109; CHECK: { 110; CHECK-NEXT: .reg .pred %p<2>; 111; CHECK-NEXT: .reg .b16 %rs<4>; 112; CHECK-NEXT: .reg .f64 %fd<5>; 113; CHECK-EMPTY: 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; 124; CHECK-NEXT: ret; 125 %c = fpext half %b to double 126 %val = call double @llvm.copysign.f64(double %a, double %c) 127 ret double %val 128} 129 130 131declare float @llvm.copysign.f32(float, float) 132declare double @llvm.copysign.f64(double, double) 133