133 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s3; 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 .b32 %r<4>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b32 %r1, [fcopysign_f_f_param_0];15; CHECK-NEXT: ld.param.b32 %r2, [fcopysign_f_f_param_1];16; CHECK-NEXT: copysign.f32 %r3, %r2, %r1;17; CHECK-NEXT: st.param.b32 [func_retval0], %r3;18; CHECK-NEXT: ret;19 %val = call float @llvm.copysign.f32(float %a, float %b)20 ret float %val21}22 23define double @fcopysign_d_d(double %a, double %b) {24; CHECK-LABEL: fcopysign_d_d(25; CHECK: {26; CHECK-NEXT: .reg .b64 %rd<4>;27; CHECK-EMPTY:28; CHECK-NEXT: // %bb.0:29; CHECK-NEXT: ld.param.b64 %rd1, [fcopysign_d_d_param_0];30; CHECK-NEXT: ld.param.b64 %rd2, [fcopysign_d_d_param_1];31; CHECK-NEXT: copysign.f64 %rd3, %rd2, %rd1;32; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;33; CHECK-NEXT: ret;34 %val = call double @llvm.copysign.f64(double %a, double %b)35 ret double %val36}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 .b32 %r<5>;43; CHECK-NEXT: .reg .b64 %rd<4>;44; CHECK-EMPTY:45; CHECK-NEXT: // %bb.0:46; CHECK-NEXT: ld.param.b32 %r1, [fcopysign_f_d_param_0];47; CHECK-NEXT: abs.f32 %r2, %r1;48; CHECK-NEXT: neg.f32 %r3, %r2;49; CHECK-NEXT: ld.param.b64 %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.ne.b64 %p1, %rd3, 0;53; CHECK-NEXT: selp.f32 %r4, %r3, %r2, %p1;54; CHECK-NEXT: st.param.b32 [func_retval0], %r4;55; CHECK-NEXT: ret;56 %c = fptrunc double %b to float57 %val = call float @llvm.copysign.f32(float %a, float %c)58 ret float %val59}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 .b32 %r<5>;67; CHECK-EMPTY:68; CHECK-NEXT: // %bb.0:69; CHECK-NEXT: ld.param.b32 %r1, [fcopysign_f_h_param_0];70; CHECK-NEXT: abs.f32 %r2, %r1;71; CHECK-NEXT: neg.f32 %r3, %r2;72; CHECK-NEXT: ld.param.b16 %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.ne.b16 %p1, %rs3, 0;76; CHECK-NEXT: selp.f32 %r4, %r3, %r2, %p1;77; CHECK-NEXT: st.param.b32 [func_retval0], %r4;78; CHECK-NEXT: ret;79 %c = fpext half %b to float80 %val = call float @llvm.copysign.f32(float %a, float %c)81 ret float %val82}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 .b64 %rd<5>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: ld.param.b64 %rd1, [fcopysign_d_f_param_0];93; CHECK-NEXT: abs.f64 %rd2, %rd1;94; CHECK-NEXT: neg.f64 %rd3, %rd2;95; CHECK-NEXT: ld.param.b32 %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.ne.b32 %p1, %r3, 0;99; CHECK-NEXT: selp.f64 %rd4, %rd3, %rd2, %p1;100; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;101; CHECK-NEXT: ret;102 %c = fpext float %b to double103 %val = call double @llvm.copysign.f64(double %a, double %c)104 ret double %val105}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 .b64 %rd<5>;113; CHECK-EMPTY:114; CHECK-NEXT: // %bb.0:115; CHECK-NEXT: ld.param.b64 %rd1, [fcopysign_d_h_param_0];116; CHECK-NEXT: abs.f64 %rd2, %rd1;117; CHECK-NEXT: neg.f64 %rd3, %rd2;118; CHECK-NEXT: ld.param.b16 %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.ne.b16 %p1, %rs3, 0;122; CHECK-NEXT: selp.f64 %rd4, %rd3, %rd2, %p1;123; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;124; CHECK-NEXT: ret;125 %c = fpext half %b to double126 %val = call double @llvm.copysign.f64(double %a, double %c)127 ret double %val128}129 130 131declare float @llvm.copysign.f32(float, float)132declare double @llvm.copysign.f64(double, double)133