421 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// Test with OCML_BASIC_ROUNDED_OPERATIONS3// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \4// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \5// RUN: -internal-isystem %S/Inputs/include \6// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \7// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \8// RUN: -D__HIPCC_RTC__ -DOCML_BASIC_ROUNDED_OPERATIONS | FileCheck %s9 10// CHECK-LABEL: @test___fadd_rd(11// CHECK-NEXT: entry:12// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3:[0-9]+]]13// CHECK-NEXT: ret float [[CALL_I]]14//15extern "C" __device__ float test___fadd_rd(float x, float y) {16 return __fadd_rd(x, y);17}18 19// CHECK-LABEL: @test___fadd_rn(20// CHECK-NEXT: entry:21// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]22// CHECK-NEXT: ret float [[CALL_I]]23//24extern "C" __device__ float test___fadd_rn(float x, float y) {25 return __fadd_rn(x, y);26}27 28// CHECK-LABEL: @test___fadd_ru(29// CHECK-NEXT: entry:30// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]31// CHECK-NEXT: ret float [[CALL_I]]32//33extern "C" __device__ float test___fadd_ru(float x, float y) {34 return __fadd_ru(x, y);35}36 37// CHECK-LABEL: @test___fadd_rz(38// CHECK-NEXT: entry:39// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]40// CHECK-NEXT: ret float [[CALL_I]]41//42extern "C" __device__ float test___fadd_rz(float x, float y) {43 return __fadd_rz(x, y);44}45 46// CHECK-LABEL: @test__fmaf_rd(47// CHECK-NEXT: entry:48// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]49// CHECK-NEXT: ret float [[CALL_I]]50//51extern "C" __device__ float test__fmaf_rd(float x, float y, float z) {52 return __fmaf_rd(x, y, z);53}54 55// CHECK-LABEL: @test__fmaf_rn(56// CHECK-NEXT: entry:57// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]58// CHECK-NEXT: ret float [[CALL_I]]59//60extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {61 return __fmaf_rn(x, y, z);62}63 64// CHECK-LABEL: @test__fmaf_ru(65// CHECK-NEXT: entry:66// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]67// CHECK-NEXT: ret float [[CALL_I]]68//69extern "C" __device__ float test__fmaf_ru(float x, float y, float z) {70 return __fmaf_ru(x, y, z);71}72 73// CHECK-LABEL: @test__fmaf_rz(74// CHECK-NEXT: entry:75// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]76// CHECK-NEXT: ret float [[CALL_I]]77//78extern "C" __device__ float test__fmaf_rz(float x, float y, float z) {79 return __fmaf_rz(x, y, z);80}81 82// CHECK-LABEL: @test___fmul_rd(83// CHECK-NEXT: entry:84// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]85// CHECK-NEXT: ret float [[CALL_I]]86//87extern "C" __device__ float test___fmul_rd(float x, float y) {88 return __fmul_rd(x, y);89}90 91// CHECK-LABEL: @test___fmul_rn(92// CHECK-NEXT: entry:93// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]94// CHECK-NEXT: ret float [[CALL_I]]95//96extern "C" __device__ float test___fmul_rn(float x, float y) {97 return __fmul_rn(x, y);98}99 100// CHECK-LABEL: @test___fmul_ru(101// CHECK-NEXT: entry:102// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]103// CHECK-NEXT: ret float [[CALL_I]]104//105extern "C" __device__ float test___fmul_ru(float x, float y) {106 return __fmul_ru(x, y);107}108 109// CHECK-LABEL: @test___fmul_rz(110// CHECK-NEXT: entry:111// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]112// CHECK-NEXT: ret float [[CALL_I]]113//114extern "C" __device__ float test___fmul_rz(float x, float y) {115 return __fmul_rz(x, y);116}117 118// CHECK-LABEL: @test___frcp_rd(119// CHECK-NEXT: entry:120// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]121// CHECK-NEXT: ret float [[CALL_I]]122//123extern "C" __device__ float test___frcp_rd(float x) {124 return __frcp_rd(x);125}126 127// CHECK-LABEL: @test___frcp_rn(128// CHECK-NEXT: entry:129// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]130// CHECK-NEXT: ret float [[CALL_I]]131//132extern "C" __device__ float test___frcp_rn(float x) {133 return __frcp_rn(x);134}135 136// CHECK-LABEL: @test___frcp_ru(137// CHECK-NEXT: entry:138// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]139// CHECK-NEXT: ret float [[CALL_I]]140//141extern "C" __device__ float test___frcp_ru(float x) {142 return __frcp_ru(x);143}144 145// CHECK-LABEL: @test___frcp_rz(146// CHECK-NEXT: entry:147// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]148// CHECK-NEXT: ret float [[CALL_I]]149//150extern "C" __device__ float test___frcp_rz(float x) {151 return __frcp_rz(x);152}153 154// CHECK-LABEL: @test___fsqrt_rd(155// CHECK-NEXT: entry:156// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR3]]157// CHECK-NEXT: ret float [[CALL_I]]158//159extern "C" __device__ float test___fsqrt_rd(float x) {160 return __fsqrt_rd(x);161}162 163// CHECK-LABEL: @test___fsqrt_rn(164// CHECK-NEXT: entry:165// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR3]]166// CHECK-NEXT: ret float [[CALL_I]]167//168extern "C" __device__ float test___fsqrt_rn(float x) {169 return __fsqrt_rn(x);170}171 172// CHECK-LABEL: @test___fsqrt_ru(173// CHECK-NEXT: entry:174// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR3]]175// CHECK-NEXT: ret float [[CALL_I]]176//177extern "C" __device__ float test___fsqrt_ru(float x) {178 return __fsqrt_ru(x);179}180 181// CHECK-LABEL: @test___fsqrt_rz(182// CHECK-NEXT: entry:183// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR3]]184// CHECK-NEXT: ret float [[CALL_I]]185//186extern "C" __device__ float test___fsqrt_rz(float x) {187 return __fsqrt_rz(x);188}189 190// CHECK-LABEL: @test___fsub_rd(191// CHECK-NEXT: entry:192// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]193// CHECK-NEXT: ret float [[CALL_I]]194//195extern "C" __device__ float test___fsub_rd(float x, float y) {196 return __fsub_rd(x, y);197}198 199// CHECK-LABEL: @test___fsub_rn(200// CHECK-NEXT: entry:201// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]202// CHECK-NEXT: ret float [[CALL_I]]203//204extern "C" __device__ float test___fsub_rn(float x, float y) {205 return __fsub_rn(x, y);206}207 208// CHECK-LABEL: @test___fsub_ru(209// CHECK-NEXT: entry:210// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]211// CHECK-NEXT: ret float [[CALL_I]]212//213extern "C" __device__ float test___fsub_ru(float x, float y) {214 return __fsub_ru(x, y);215}216 217// CHECK-LABEL: @test___fsub_rz(218// CHECK-NEXT: entry:219// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]220// CHECK-NEXT: ret float [[CALL_I]]221//222extern "C" __device__ float test___fsub_rz(float x, float y) {223 return __fsub_rz(x, y);224}225 226// CHECK-LABEL: @test___dadd_rd(227// CHECK-NEXT: entry:228// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]229// CHECK-NEXT: ret double [[CALL_I]]230//231extern "C" __device__ double test___dadd_rd(double x, double y) {232 return __dadd_rd(x, y);233}234 235// CHECK-LABEL: @test___dadd_rn(236// CHECK-NEXT: entry:237// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]238// CHECK-NEXT: ret double [[CALL_I]]239//240extern "C" __device__ double test___dadd_rn(double x, double y) {241 return __dadd_rn(x, y);242}243 244// CHECK-LABEL: @test___dadd_ru(245// CHECK-NEXT: entry:246// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]247// CHECK-NEXT: ret double [[CALL_I]]248//249extern "C" __device__ double test___dadd_ru(double x, double y) {250 return __dadd_ru(x, y);251}252 253// CHECK-LABEL: @test___dadd_rz(254// CHECK-NEXT: entry:255// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]256// CHECK-NEXT: ret double [[CALL_I]]257//258extern "C" __device__ double test___dadd_rz(double x, double y) {259 return __dadd_rz(x, y);260}261 262// CHECK-LABEL: @test___dmul_rd(263// CHECK-NEXT: entry:264// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]265// CHECK-NEXT: ret double [[CALL_I]]266//267extern "C" __device__ double test___dmul_rd(double x, double y) {268 return __dmul_rd(x, y);269}270 271// CHECK-LABEL: @test___dmul_rn(272// CHECK-NEXT: entry:273// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]274// CHECK-NEXT: ret double [[CALL_I]]275//276extern "C" __device__ double test___dmul_rn(double x, double y) {277 return __dmul_rn(x, y);278}279 280// CHECK-LABEL: @test___dmul_ru(281// CHECK-NEXT: entry:282// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]283// CHECK-NEXT: ret double [[CALL_I]]284//285extern "C" __device__ double test___dmul_ru(double x, double y) {286 return __dmul_ru(x, y);287}288 289// CHECK-LABEL: @test___dmul_rz(290// CHECK-NEXT: entry:291// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]292// CHECK-NEXT: ret double [[CALL_I]]293//294extern "C" __device__ double test___dmul_rz(double x, double y) {295 return __dmul_rz(x, y);296}297 298// CHECK-LABEL: @test___drcp_rd(299// CHECK-NEXT: entry:300// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double301// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]302// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float303// CHECK-NEXT: ret float [[CONV1]]304//305extern "C" __device__ float test___drcp_rd(float x) {306 return __drcp_rd(x);307}308 309// CHECK-LABEL: @test___drcp_rn(310// CHECK-NEXT: entry:311// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double312// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]313// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float314// CHECK-NEXT: ret float [[CONV1]]315//316extern "C" __device__ float test___drcp_rn(float x) {317 return __drcp_rn(x);318}319 320// CHECK-LABEL: @test___drcp_ru(321// CHECK-NEXT: entry:322// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double323// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]324// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float325// CHECK-NEXT: ret float [[CONV1]]326//327extern "C" __device__ float test___drcp_ru(float x) {328 return __drcp_ru(x);329}330 331// CHECK-LABEL: @test___drcp_rz(332// CHECK-NEXT: entry:333// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double334// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]335// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float336// CHECK-NEXT: ret float [[CONV1]]337//338extern "C" __device__ float test___drcp_rz(float x) {339 return __drcp_rz(x);340}341 342// CHECK-LABEL: @test___dsqrt_rd(343// CHECK-NEXT: entry:344// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double345// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR3]]346// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float347// CHECK-NEXT: ret float [[CONV1]]348//349extern "C" __device__ float test___dsqrt_rd(float x) {350 return __dsqrt_rd(x);351}352 353// CHECK-LABEL: @test___dsqrt_rn(354// CHECK-NEXT: entry:355// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double356// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR3]]357// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float358// CHECK-NEXT: ret float [[CONV1]]359//360extern "C" __device__ float test___dsqrt_rn(float x) {361 return __dsqrt_rn(x);362}363 364// CHECK-LABEL: @test___dsqrt_ru(365// CHECK-NEXT: entry:366// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double367// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR3]]368// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float369// CHECK-NEXT: ret float [[CONV1]]370//371extern "C" __device__ float test___dsqrt_ru(float x) {372 return __dsqrt_ru(x);373}374 375// CHECK-LABEL: @test___dsqrt_rz(376// CHECK-NEXT: entry:377// CHECK-NEXT: [[CONV:%.*]] = fpext contract float [[X:%.*]] to double378// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR3]]379// CHECK-NEXT: [[CONV1:%.*]] = fptrunc contract double [[CALL_I]] to float380// CHECK-NEXT: ret float [[CONV1]]381//382extern "C" __device__ float test___dsqrt_rz(float x) {383 return __dsqrt_rz(x);384}385 386// CHECK-LABEL: @test__fma_rd(387// CHECK-NEXT: entry:388// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]389// CHECK-NEXT: ret double [[CALL_I]]390//391extern "C" __device__ double test__fma_rd(double x, double y, double z) {392 return __fma_rd(x, y, z);393}394 395// CHECK-LABEL: @test__fma_rn(396// CHECK-NEXT: entry:397// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]398// CHECK-NEXT: ret double [[CALL_I]]399//400extern "C" __device__ double test__fma_rn(double x, double y, double z) {401 return __fma_rn(x, y, z);402}403 404// CHECK-LABEL: @test__fma_ru(405// CHECK-NEXT: entry:406// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]407// CHECK-NEXT: ret double [[CALL_I]]408//409extern "C" __device__ double test__fma_ru(double x, double y, double z) {410 return __fma_ru(x, y, z);411}412 413// CHECK-LABEL: @test__fma_rz(414// CHECK-NEXT: entry:415// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]416// CHECK-NEXT: ret double [[CALL_I]]417//418extern "C" __device__ double test__fma_rz(double x, double y, double z) {419 return __fma_rz(x, y, z);420}421