This is an archive of the discontinued LLVM Phabricator instance.

Differential D152858

OpenMP: Use generated checks and pragma declare target
ClosedPublic

Authored by arsenm on Jun 13 2023, 2:33 PM.

Details

Reviewers
JonChesterfield
doru1004
ronlieb
ddpagan
jdoerfert
mhalk
Group Reviewers
Restricted Project

Diff Detail

Event Timeline

arsenm created this revision.Jun 13 2023, 2:33 PM
Herald added a project: Restricted Project. · View Herald TranscriptJun 13 2023, 2:33 PM
Herald added subscribers: sunshaoce, guansong, yaxunl, jvesely. · View Herald Transcript
arsenm requested review of this revision.Jun 13 2023, 2:33 PM
Herald added subscribers: jplehr, sstefan1, wdng. · View Herald TranscriptJun 13 2023, 2:33 PM
Harbormaster completed remote builds in B238617: Diff 531072.Jun 13 2023, 3:55 PM
arsenm added a reviewer: mhalk.Jun 19 2023, 4:49 AM
arsenm added a comment.Jun 22 2023, 9:45 AM

ping

jdoerfert accepted this revision.Jun 26 2023, 11:24 AM
This revision is now accepted and ready to land.Jun 26 2023, 11:24 AM
arsenm closed this revision.Jun 26 2023, 11:45 AM

53d28b2a71063b974ec058ced614d1c95d5584fe

Revision Contents

PathSize
clang/
test/
Headers/
318 lines

Diff 531072

clang/test/Headers/amdgcn_openmp_device_math.c

// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc // RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK-C,CHECK // RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK,CHECK-C
// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc // RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK-CPP,CHECK // RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK,CHECK-CPP
#ifdef __cplusplus #ifdef __cplusplus
#include <cmath> #include <cmath>
#else #else
#include <math.h> #include <math.h>
#endif #endif
#pragma omp begin declare target
// CHECK-C-LABEL: @test_math_f64(
// CHECK-C-NEXT: entry:
// CHECK-C-NEXT: [[RETVAL_I8:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I9:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[RETVAL_I3:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I4:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[RETVAL_I:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[X_ADDR:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[L1:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[L2:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[L3:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-C-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-C-NEXT: [[L2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L2]] to ptr
// CHECK-C-NEXT: [[L3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L3]] to ptr
// CHECK-C-NEXT: store double [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK-C-NEXT: [[TMP0:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-C-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-C-NEXT: store double [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-C-NEXT: [[TMP1:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-C-NEXT: [[CALL_I:%.*]] = call double @__ocml_sin_f64(double noundef [[TMP1]]) #[[ATTR3:[0-9]+]]
// CHECK-C-NEXT: store double [[CALL_I]], ptr [[L1_ASCAST]], align 8
// CHECK-C-NEXT: [[TMP2:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-C-NEXT: [[RETVAL_ASCAST_I5:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I3]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I4]] to ptr
// CHECK-C-NEXT: store double [[TMP2]], ptr [[__X_ADDR_ASCAST_I6]], align 8
// CHECK-C-NEXT: [[TMP3:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I6]], align 8
// CHECK-C-NEXT: [[CALL_I7:%.*]] = call double @__ocml_cos_f64(double noundef [[TMP3]]) #[[ATTR3]]
// CHECK-C-NEXT: store double [[CALL_I7]], ptr [[L2_ASCAST]], align 8
// CHECK-C-NEXT: [[TMP4:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-C-NEXT: [[RETVAL_ASCAST_I10:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I8]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I9]] to ptr
// CHECK-C-NEXT: store double [[TMP4]], ptr [[__X_ADDR_ASCAST_I11]], align 8
// CHECK-C-NEXT: [[TMP5:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I11]], align 8
// CHECK-C-NEXT: [[TMP6:%.*]] = call double @llvm.fabs.f64(double [[TMP5]])
// CHECK-C-NEXT: store double [[TMP6]], ptr [[L3_ASCAST]], align 8
// CHECK-C-NEXT: ret void
//
// CHECK-CPP-LABEL: @_Z13test_math_f64d(
// CHECK-CPP-NEXT: entry:
// CHECK-CPP-NEXT: [[RETVAL_I8:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I9:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I3:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I4:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[X_ADDR:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[L1:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[L2:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[L3:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-CPP-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-CPP-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-CPP-NEXT: [[L2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L2]] to ptr
// CHECK-CPP-NEXT: [[L3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L3]] to ptr
// CHECK-CPP-NEXT: store double [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK-CPP-NEXT: [[TMP0:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-CPP-NEXT: store double [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-CPP-NEXT: [[TMP1:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call double @__ocml_sin_f64(double noundef [[TMP1]]) #[[ATTR3:[0-9]+]]
// CHECK-CPP-NEXT: store double [[CALL_I]], ptr [[L1_ASCAST]], align 8
// CHECK-CPP-NEXT: [[TMP2:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I5:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I3]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I4]] to ptr
// CHECK-CPP-NEXT: store double [[TMP2]], ptr [[__X_ADDR_ASCAST_I6]], align 8
// CHECK-CPP-NEXT: [[TMP3:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I6]], align 8
// CHECK-CPP-NEXT: [[CALL_I7:%.*]] = call double @__ocml_cos_f64(double noundef [[TMP3]]) #[[ATTR3]]
// CHECK-CPP-NEXT: store double [[CALL_I7]], ptr [[L2_ASCAST]], align 8
// CHECK-CPP-NEXT: [[TMP4:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I10:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I8]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I9]] to ptr
// CHECK-CPP-NEXT: store double [[TMP4]], ptr [[__X_ADDR_ASCAST_I11]], align 8
// CHECK-CPP-NEXT: [[TMP5:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I11]], align 8
// CHECK-CPP-NEXT: [[TMP6:%.*]] = call double @llvm.fabs.f64(double [[TMP5]])
// CHECK-CPP-NEXT: store double [[TMP6]], ptr [[L3_ASCAST]], align 8
// CHECK-CPP-NEXT: ret void
//
void test_math_f64(double x) { void test_math_f64(double x) {
// CHECK-LABEL: define {{.*}}test_math_f64
#pragma omp target
{
// CHECK: call double @__ocml_sin_f64
double l1 = sin(x); double l1 = sin(x);
// CHECK: call double @__ocml_cos_f64
double l2 = cos(x); double l2 = cos(x);
// CHECK: call double @llvm.fabs.f64
double l3 = fabs(x); double l3 = fabs(x);
} }
}
// CHECK-C-LABEL: @test_math_f32(
// CHECK-C-NEXT: entry:
// CHECK-C-NEXT: [[RETVAL_I13:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I14:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[RETVAL_I8:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I9:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[RETVAL_I:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I:%.*]] = alloca double, align 8, addrspace(5)
// CHECK-C-NEXT: [[X_ADDR:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[L1:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[L2:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[L3:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-C-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-C-NEXT: [[L2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L2]] to ptr
// CHECK-C-NEXT: [[L3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L3]] to ptr
// CHECK-C-NEXT: store float [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[TMP0:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[CONV:%.*]] = fpext float [[TMP0]] to double
// CHECK-C-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-C-NEXT: store double [[CONV]], ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-C-NEXT: [[TMP1:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-C-NEXT: [[CALL_I:%.*]] = call double @__ocml_sin_f64(double noundef [[TMP1]]) #[[ATTR3]]
// CHECK-C-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
// CHECK-C-NEXT: store float [[CONV1]], ptr [[L1_ASCAST]], align 4
// CHECK-C-NEXT: [[TMP2:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[CONV2:%.*]] = fpext float [[TMP2]] to double
// CHECK-C-NEXT: [[RETVAL_ASCAST_I10:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I8]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I9]] to ptr
// CHECK-C-NEXT: store double [[CONV2]], ptr [[__X_ADDR_ASCAST_I11]], align 8
// CHECK-C-NEXT: [[TMP3:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I11]], align 8
// CHECK-C-NEXT: [[CALL_I12:%.*]] = call double @__ocml_cos_f64(double noundef [[TMP3]]) #[[ATTR3]]
// CHECK-C-NEXT: [[CONV4:%.*]] = fptrunc double [[CALL_I12]] to float
// CHECK-C-NEXT: store float [[CONV4]], ptr [[L2_ASCAST]], align 4
// CHECK-C-NEXT: [[TMP4:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[CONV5:%.*]] = fpext float [[TMP4]] to double
// CHECK-C-NEXT: [[RETVAL_ASCAST_I15:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I13]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I16:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I14]] to ptr
// CHECK-C-NEXT: store double [[CONV5]], ptr [[__X_ADDR_ASCAST_I16]], align 8
// CHECK-C-NEXT: [[TMP5:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I16]], align 8
// CHECK-C-NEXT: [[TMP6:%.*]] = call double @llvm.fabs.f64(double [[TMP5]])
// CHECK-C-NEXT: [[CONV7:%.*]] = fptrunc double [[TMP6]] to float
// CHECK-C-NEXT: store float [[CONV7]], ptr [[L3_ASCAST]], align 4
// CHECK-C-NEXT: ret void
//
// CHECK-CPP-LABEL: @_Z13test_math_f32f(
// CHECK-CPP-NEXT: entry:
// CHECK-CPP-NEXT: [[RETVAL_I22:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I23:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I18:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I19:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I13:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I14:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I8:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I9:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I3:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I4:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[X_ADDR:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[L1:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[L2:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[L3:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-CPP-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-CPP-NEXT: [[L2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L2]] to ptr
// CHECK-CPP-NEXT: [[L3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L3]] to ptr
// CHECK-CPP-NEXT: store float [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[TMP0:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-CPP-NEXT: store float [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-CPP-NEXT: [[TMP1:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I15:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I13]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I16:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I14]] to ptr
// CHECK-CPP-NEXT: store float [[TMP1]], ptr [[__X_ADDR_ASCAST_I16]], align 4
// CHECK-CPP-NEXT: [[TMP2:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I16]], align 4
// CHECK-CPP-NEXT: [[CALL_I17:%.*]] = call float @__ocml_sin_f32(float noundef [[TMP2]]) #[[ATTR3]]
// CHECK-CPP-NEXT: store float [[CALL_I17]], ptr [[L1_ASCAST]], align 4
// CHECK-CPP-NEXT: [[TMP3:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I5:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I3]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I4]] to ptr
// CHECK-CPP-NEXT: store float [[TMP3]], ptr [[__X_ADDR_ASCAST_I6]], align 4
// CHECK-CPP-NEXT: [[TMP4:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I6]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I20:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I18]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I21:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I19]] to ptr
// CHECK-CPP-NEXT: store float [[TMP4]], ptr [[__X_ADDR_ASCAST_I21]], align 4
// CHECK-CPP-NEXT: [[TMP5:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I21]], align 4
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call float @__ocml_cos_f32(float noundef [[TMP5]]) #[[ATTR3]]
// CHECK-CPP-NEXT: store float [[CALL_I]], ptr [[L2_ASCAST]], align 4
// CHECK-CPP-NEXT: [[TMP6:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I10:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I8]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I9]] to ptr
// CHECK-CPP-NEXT: store float [[TMP6]], ptr [[__X_ADDR_ASCAST_I11]], align 4
// CHECK-CPP-NEXT: [[TMP7:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I11]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I24:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I22]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I25:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I23]] to ptr
// CHECK-CPP-NEXT: store float [[TMP7]], ptr [[__X_ADDR_ASCAST_I25]], align 4
// CHECK-CPP-NEXT: [[TMP8:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I25]], align 4
// CHECK-CPP-NEXT: [[TMP9:%.*]] = call float @llvm.fabs.f32(float [[TMP8]])
// CHECK-CPP-NEXT: store float [[TMP9]], ptr [[L3_ASCAST]], align 4
// CHECK-CPP-NEXT: ret void
//
void test_math_f32(float x) { void test_math_f32(float x) {
// CHECK-LABEL: define {{.*}}test_math_f32
#pragma omp target
{
// CHECK-C: call double @__ocml_sin_f64
// CHECK-CPP: call float @__ocml_sin_f32
float l1 = sin(x); float l1 = sin(x);
// CHECK-C: call double @__ocml_cos_f64
// CHECK-CPP: call float @__ocml_cos_f32
float l2 = cos(x); float l2 = cos(x);
// CHECK-C: call double @llvm.fabs.f64
// CHECK-CPP: call float @llvm.fabs.f32
float l3 = fabs(x); float l3 = fabs(x);
} }
}
// CHECK-C-LABEL: @test_math_f32_suffix(
// CHECK-C-NEXT: entry:
// CHECK-C-NEXT: [[RETVAL_I8:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I9:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[RETVAL_I3:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I4:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[RETVAL_I:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[__X_ADDR_I:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[X_ADDR:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[L1:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[L2:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[L3:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-C-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-C-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-C-NEXT: [[L2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L2]] to ptr
// CHECK-C-NEXT: [[L3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L3]] to ptr
// CHECK-C-NEXT: store float [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[TMP0:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-C-NEXT: store float [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-C-NEXT: [[TMP1:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-C-NEXT: [[CALL_I:%.*]] = call float @__ocml_sin_f32(float noundef [[TMP1]]) #[[ATTR3]]
// CHECK-C-NEXT: store float [[CALL_I]], ptr [[L1_ASCAST]], align 4
// CHECK-C-NEXT: [[TMP2:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[RETVAL_ASCAST_I5:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I3]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I4]] to ptr
// CHECK-C-NEXT: store float [[TMP2]], ptr [[__X_ADDR_ASCAST_I6]], align 4
// CHECK-C-NEXT: [[TMP3:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I6]], align 4
// CHECK-C-NEXT: [[CALL_I7:%.*]] = call float @__ocml_cos_f32(float noundef [[TMP3]]) #[[ATTR3]]
// CHECK-C-NEXT: store float [[CALL_I7]], ptr [[L2_ASCAST]], align 4
// CHECK-C-NEXT: [[TMP4:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-C-NEXT: [[RETVAL_ASCAST_I10:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I8]] to ptr
// CHECK-C-NEXT: [[__X_ADDR_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I9]] to ptr
// CHECK-C-NEXT: store float [[TMP4]], ptr [[__X_ADDR_ASCAST_I11]], align 4
// CHECK-C-NEXT: [[TMP5:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I11]], align 4
// CHECK-C-NEXT: [[TMP6:%.*]] = call float @llvm.fabs.f32(float [[TMP5]])
// CHECK-C-NEXT: store float [[TMP6]], ptr [[L3_ASCAST]], align 4
// CHECK-C-NEXT: ret void
//
// CHECK-CPP-LABEL: @_Z20test_math_f32_suffixf(
// CHECK-CPP-NEXT: entry:
// CHECK-CPP-NEXT: [[RETVAL_I8:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I9:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I3:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I4:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[RETVAL_I:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[__X_ADDR_I:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[X_ADDR:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[L1:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[L2:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[L3:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-CPP-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-CPP-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-CPP-NEXT: [[L2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L2]] to ptr
// CHECK-CPP-NEXT: [[L3_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L3]] to ptr
// CHECK-CPP-NEXT: store float [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[TMP0:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-CPP-NEXT: store float [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-CPP-NEXT: [[TMP1:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call float @__ocml_sin_f32(float noundef [[TMP1]]) #[[ATTR3]]
// CHECK-CPP-NEXT: store float [[CALL_I]], ptr [[L1_ASCAST]], align 4
// CHECK-CPP-NEXT: [[TMP2:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I5:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I3]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I4]] to ptr
// CHECK-CPP-NEXT: store float [[TMP2]], ptr [[__X_ADDR_ASCAST_I6]], align 4
// CHECK-CPP-NEXT: [[TMP3:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I6]], align 4
// CHECK-CPP-NEXT: [[CALL_I7:%.*]] = call float @__ocml_cos_f32(float noundef [[TMP3]]) #[[ATTR3]]
// CHECK-CPP-NEXT: store float [[CALL_I7]], ptr [[L2_ASCAST]], align 4
// CHECK-CPP-NEXT: [[TMP4:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I10:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I8]] to ptr
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I9]] to ptr
// CHECK-CPP-NEXT: store float [[TMP4]], ptr [[__X_ADDR_ASCAST_I11]], align 4
// CHECK-CPP-NEXT: [[TMP5:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I11]], align 4
// CHECK-CPP-NEXT: [[TMP6:%.*]] = call float @llvm.fabs.f32(float [[TMP5]])
// CHECK-CPP-NEXT: store float [[TMP6]], ptr [[L3_ASCAST]], align 4
// CHECK-CPP-NEXT: ret void
//
void test_math_f32_suffix(float x) { void test_math_f32_suffix(float x) {
// CHECK-LABEL: define {{.*}}test_math_f32_suffix
#pragma omp target
{
// CHECK: call float @__ocml_sin_f32
float l1 = sinf(x); float l1 = sinf(x);
// CHECK: call float @__ocml_cos_f32
float l2 = cosf(x); float l2 = cosf(x);
// CHECK: call float @llvm.fabs.f32
float l3 = fabsf(x); float l3 = fabsf(x);
} }
}
#pragma omp end declare target
//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
// CHECK: {{.*}}