test/Headers/__clang_hip_math_ocml_rounded_ops.hip

// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// Test with OCML_BASIC_ROUNDED_OPERATIONS
// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h      \
// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN:   -internal-isystem %S/Inputs/include \
// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
// RUN:   -D__HIPCC_RTC__ -DOCML_BASIC_ROUNDED_OPERATIONS | FileCheck %s

// CHECK-LABEL: @test___fadd_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4:[0-9]+]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fadd_rd(float x, float y) {
  return __fadd_rd(x, y);
}

// CHECK-LABEL: @test___fadd_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fadd_rn(float x, float y) {
  return __fadd_rn(x, y);
}

// CHECK-LABEL: @test___fadd_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fadd_ru(float x, float y) {
  return __fadd_ru(x, y);
}

// CHECK-LABEL: @test___fadd_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fadd_rz(float x, float y) {
  return __fadd_rz(x, y);
}

// CHECK-LABEL: @test__fmaf_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test__fmaf_rd(float x, float y, float z) {
  return __fmaf_rd(x, y, z);
}

// CHECK-LABEL: @test__fmaf_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
  return __fmaf_rn(x, y, z);
}

// CHECK-LABEL: @test__fmaf_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test__fmaf_ru(float x, float y, float z) {
  return __fmaf_ru(x, y, z);
}

// CHECK-LABEL: @test__fmaf_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test__fmaf_rz(float x, float y, float z) {
  return __fmaf_rz(x, y, z);
}

// CHECK-LABEL: @test___fmul_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fmul_rd(float x, float y) {
  return __fmul_rd(x, y);
}

// CHECK-LABEL: @test___fmul_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fmul_rn(float x, float y) {
  return __fmul_rn(x, y);
}

// CHECK-LABEL: @test___fmul_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fmul_ru(float x, float y) {
  return __fmul_ru(x, y);
}

// CHECK-LABEL: @test___fmul_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fmul_rz(float x, float y) {
  return __fmul_rz(x, y);
}

// CHECK-LABEL: @test___frcp_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___frcp_rd(float x) {
  return __frcp_rd(x);
}

// CHECK-LABEL: @test___frcp_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___frcp_rn(float x) {
  return __frcp_rn(x);
}

// CHECK-LABEL: @test___frcp_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___frcp_ru(float x) {
  return __frcp_ru(x);
}

// CHECK-LABEL: @test___frcp_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___frcp_rz(float x) {
  return __frcp_rz(x);
}

// CHECK-LABEL: @test___fsqrt_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsqrt_rd(float x) {
  return __fsqrt_rd(x);
}

// CHECK-LABEL: @test___fsqrt_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsqrt_rn(float x) {
  return __fsqrt_rn(x);
}

// CHECK-LABEL: @test___fsqrt_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsqrt_ru(float x) {
  return __fsqrt_ru(x);
}

// CHECK-LABEL: @test___fsqrt_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsqrt_rz(float x) {
  return __fsqrt_rz(x);
}

// CHECK-LABEL: @test___fsub_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsub_rd(float x, float y) {
  return __fsub_rd(x, y);
}

// CHECK-LABEL: @test___fsub_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsub_rn(float x, float y) {
  return __fsub_rn(x, y);
}

// CHECK-LABEL: @test___fsub_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsub_ru(float x, float y) {
  return __fsub_ru(x, y);
}

// CHECK-LABEL: @test___fsub_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret float [[CALL_I]]
//
extern "C" __device__ float test___fsub_rz(float x, float y) {
  return __fsub_rz(x, y);
}

// CHECK-LABEL: @test___dadd_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dadd_rd(double x, double y) {
  return __dadd_rd(x, y);
}

// CHECK-LABEL: @test___dadd_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dadd_rn(double x, double y) {
  return __dadd_rn(x, y);
}

// CHECK-LABEL: @test___dadd_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dadd_ru(double x, double y) {
  return __dadd_ru(x, y);
}

// CHECK-LABEL: @test___dadd_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dadd_rz(double x, double y) {
  return __dadd_rz(x, y);
}

// CHECK-LABEL: @test___dmul_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dmul_rd(double x, double y) {
  return __dmul_rd(x, y);
}

// CHECK-LABEL: @test___dmul_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dmul_rn(double x, double y) {
  return __dmul_rn(x, y);
}

// CHECK-LABEL: @test___dmul_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dmul_ru(double x, double y) {
  return __dmul_ru(x, y);
}

// CHECK-LABEL: @test___dmul_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test___dmul_rz(double x, double y) {
  return __dmul_rz(x, y);
}

// CHECK-LABEL: @test___drcp_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___drcp_rd(float x) {
  return __drcp_rd(x);
}

// CHECK-LABEL: @test___drcp_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___drcp_rn(float x) {
  return __drcp_rn(x);
}

// CHECK-LABEL: @test___drcp_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___drcp_ru(float x) {
  return __drcp_ru(x);
}

// CHECK-LABEL: @test___drcp_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___drcp_rz(float x) {
  return __drcp_rz(x);
}

// CHECK-LABEL: @test___dsqrt_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___dsqrt_rd(float x) {
  return __dsqrt_rd(x);
}

// CHECK-LABEL: @test___dsqrt_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___dsqrt_rn(float x) {
  return __dsqrt_rn(x);
}

// CHECK-LABEL: @test___dsqrt_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___dsqrt_ru(float x) {
  return __dsqrt_ru(x);
}

// CHECK-LABEL: @test___dsqrt_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR4]]
// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-NEXT:    ret float [[CONV1]]
//
extern "C" __device__ float test___dsqrt_rz(float x) {
  return __dsqrt_rz(x);
}

// CHECK-LABEL: @test__fma_rd(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test__fma_rd(double x, double y, double z) {
  return __fma_rd(x, y, z);
}

// CHECK-LABEL: @test__fma_rn(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test__fma_rn(double x, double y, double z) {
  return __fma_rn(x, y, z);
}

// CHECK-LABEL: @test__fma_ru(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test__fma_ru(double x, double y, double z) {
  return __fma_ru(x, y, z);
}

// CHECK-LABEL: @test__fma_rz(
// CHECK-NEXT:  entry:
// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
// CHECK-NEXT:    ret double [[CALL_I]]
//
extern "C" __device__ double test__fma_rz(double x, double y, double z) {
  return __fma_rz(x, y, z);
}