145 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 22// REQUIRES: amdgpu-registered-target, x86-registered-target3 4// fp-contract, -no-enable-noundef-analysis and visibility are to just get the5// same output for openmp and hip.6 7// RUN: %clang_cc1 -x c++ -fopenmp -fvisibility=default -ffp-contract=off -no-enable-noundef-analysis \8// RUN: -internal-isystem %S/../../lib/Headers/openmp_wrappers \9// RUN: -include __clang_openmp_device_functions.h \10// RUN: -internal-isystem %S/../../lib/Headers/openmp_wrappers \11// RUN: -internal-isystem %S/Inputs/include \12// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -verify \13// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm -fopenmp-is-target-device \14// RUN: -o - %s | FileCheck --check-prefixes=CHECK,OPENMP,OPENMP-CPP %s15 16// RUN: %clang_cc1 -x c -fopenmp -fvisibility=default -ffp-contract=off -no-enable-noundef-analysis \17// RUN: -internal-isystem %S/../../lib/Headers/openmp_wrappers \18// RUN: -include __clang_openmp_device_functions.h \19// RUN: -internal-isystem %S/../../lib/Headers/openmp_wrappers \20// RUN: -internal-isystem %S/Inputs/include \21// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -verify \22// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm -fopenmp-is-target-device \23// RUN: -o - %s | FileCheck --check-prefixes=CHECK,OPENMP,OPENMP-C %s24 25// RUN: %clang_cc1 -x hip -fvisibility=default -ffp-contract=off -no-enable-noundef-analysis \26// RUN: -include __clang_hip_runtime_wrapper.h \27// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \28// RUN: -internal-isystem %S/Inputs/include \29// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -verify \30// RUN: -emit-llvm -fcuda-is-device -o - \31// RUN: -D__HIPCC_RTC__ %s | FileCheck -check-prefixes=CHECK,HIP %s32 33// expected-no-diagnostics34 35#ifdef __cplusplus36#include <cmath>37#else38#include <math.h>39typedef _Bool bool;40#endif41 42#ifdef _OPENMP43#define __device__44#endif45 46// static and overloadable to get the same function annotations between C, C++ and HIP.47#define TEST_FUNC_ATTRS static __device__ __attribute__((used,overloadable))48 49#ifdef _OPENMP50#pragma omp begin declare target51#endif52 53// This function is a hack to get the same IR out of HIP and OpenMP. The HIP54// headers declare __cxa_* functions with these attributes, such that the55// attribute groups are different. update_cc_test_checks isn't smart enough to56// strip attributes from the checks, or semantically compare them.57__attribute__((__visibility__("default")))58__attribute__((weak))59__attribute__((noreturn))60__attribute__((overloadable))61// OPENMP-LABEL: define weak hidden void @_Z20attribute_check_hackv62// OPENMP-SAME: () #[[ATTR0:[0-9]+]] {63// OPENMP-NEXT: entry:64// OPENMP-NEXT: call void @llvm.trap()65// OPENMP-NEXT: unreachable66//67void attribute_check_hack(void) {68 __builtin_trap();69}70 71// CHECK-LABEL: define internal float @_ZL18test_ockl_acos_f32f72// CHECK-SAME: (float [[SRC:%.*]]) #[[ATTR2:[0-9]+]] {73// CHECK-NEXT: entry:74// CHECK-NEXT: [[RETVAL:%.*]] = alloca float, align 4, addrspace(5)75// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)76// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr77// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr78// CHECK-NEXT: store float [[SRC]], ptr [[SRC_ADDR_ASCAST]], align 479// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 480// CHECK-NEXT: [[CALL:%.*]] = call float @__ocml_acos_f32(float [[TMP0]]) #[[ATTR4:[0-9]+]]81// CHECK-NEXT: ret float [[CALL]]82//83TEST_FUNC_ATTRS float test_ockl_acos_f32(float src) {84 return __ocml_acos_f32(src);85}86 87// CHECK-LABEL: define internal float @_ZL15test_ockl_fdot2Dv2_DF16_S_fbi88// CHECK-SAME: (<2 x half> [[A:%.*]], <2 x half> [[B:%.*]], float [[C:%.*]], i1 zeroext [[S:%.*]], i32 [[S_INT:%.*]]) #[[ATTR2]] {89// CHECK-NEXT: entry:90// CHECK-NEXT: [[RETVAL:%.*]] = alloca float, align 4, addrspace(5)91// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <2 x half>, align 4, addrspace(5)92// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <2 x half>, align 4, addrspace(5)93// CHECK-NEXT: [[C_ADDR:%.*]] = alloca float, align 4, addrspace(5)94// CHECK-NEXT: [[S_ADDR:%.*]] = alloca i8, align 1, addrspace(5)95// CHECK-NEXT: [[S_INT_ADDR:%.*]] = alloca i32, align 4, addrspace(5)96// CHECK-NEXT: [[X:%.*]] = alloca float, align 4, addrspace(5)97// CHECK-NEXT: [[Y:%.*]] = alloca float, align 4, addrspace(5)98// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr99// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr100// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr101// CHECK-NEXT: [[C_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[C_ADDR]] to ptr102// CHECK-NEXT: [[S_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S_ADDR]] to ptr103// CHECK-NEXT: [[S_INT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S_INT_ADDR]] to ptr104// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr105// CHECK-NEXT: [[Y_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[Y]] to ptr106// CHECK-NEXT: store <2 x half> [[A]], ptr [[A_ADDR_ASCAST]], align 4107// CHECK-NEXT: store <2 x half> [[B]], ptr [[B_ADDR_ASCAST]], align 4108// CHECK-NEXT: store float [[C]], ptr [[C_ADDR_ASCAST]], align 4109// CHECK-NEXT: [[FROMBOOL:%.*]] = zext i1 [[S]] to i8110// CHECK-NEXT: store i8 [[FROMBOOL]], ptr [[S_ADDR_ASCAST]], align 1111// CHECK-NEXT: store i32 [[S_INT]], ptr [[S_INT_ADDR_ASCAST]], align 4112// CHECK-NEXT: [[TMP0:%.*]] = load <2 x half>, ptr [[A_ADDR_ASCAST]], align 4113// CHECK-NEXT: [[TMP1:%.*]] = load <2 x half>, ptr [[B_ADDR_ASCAST]], align 4114// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[C_ADDR_ASCAST]], align 4115// CHECK-NEXT: [[TMP3:%.*]] = load i8, ptr [[S_ADDR_ASCAST]], align 1116// CHECK-NEXT: [[TOBOOL:%.*]] = trunc i8 [[TMP3]] to i1117// CHECK-NEXT: [[CALL:%.*]] = call float @__ockl_fdot2(<2 x half> [[TMP0]], <2 x half> [[TMP1]], float [[TMP2]], i1 zeroext [[TOBOOL]]) #[[ATTR4]]118// CHECK-NEXT: store float [[CALL]], ptr [[X_ASCAST]], align 4119// CHECK-NEXT: [[TMP4:%.*]] = load <2 x half>, ptr [[A_ADDR_ASCAST]], align 4120// CHECK-NEXT: [[TMP5:%.*]] = load <2 x half>, ptr [[B_ADDR_ASCAST]], align 4121// CHECK-NEXT: [[TMP6:%.*]] = load float, ptr [[C_ADDR_ASCAST]], align 4122// CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr [[S_INT_ADDR_ASCAST]], align 4123// CHECK-NEXT: [[TOBOOL1:%.*]] = icmp ne i32 [[TMP7]], 0124// CHECK-NEXT: [[CALL2:%.*]] = call float @__ockl_fdot2(<2 x half> [[TMP4]], <2 x half> [[TMP5]], float [[TMP6]], i1 zeroext [[TOBOOL1]]) #[[ATTR4]]125// CHECK-NEXT: store float [[CALL2]], ptr [[Y_ASCAST]], align 4126// CHECK-NEXT: [[TMP8:%.*]] = load float, ptr [[X_ASCAST]], align 4127// CHECK-NEXT: [[TMP9:%.*]] = load float, ptr [[Y_ASCAST]], align 4128// CHECK-NEXT: [[ADD:%.*]] = fadd float [[TMP8]], [[TMP9]]129// CHECK-NEXT: ret float [[ADD]]130//131TEST_FUNC_ATTRS float test_ockl_fdot2(__2f16 a, __2f16 b, float c, bool s, int s_int) {132 float x = __ockl_fdot2(a, b, c, s);133 float y = __ockl_fdot2(a, b, c, s_int);134 return x + y;135}136 137 138#ifdef _OPENMP139#pragma omp end declare target140#endif141//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:142// HIP: {{.*}}143// OPENMP-C: {{.*}}144// OPENMP-CPP: {{.*}}145