1 // RUN: %clang_cc1 %s -triple x86_64-linux-unknown -emit-llvm -o - \
2 // RUN: | FileCheck -check-prefix=HOST %s
3 // RUN: %clang_cc1 %s -fcuda-is-device \
4 // RUN: -emit-llvm -o - -triple nvptx64 \
5 // RUN: -aux-triple x86_64-unknown-linux-gnu | FileCheck \
6 // RUN: -check-prefix=DEV %s
7
8 #include "Inputs/cuda.h"
9
10 // Check host/device-based overloding resolution in global variable initializer.
pow(double,double)11 double pow(double, double) { return 1.0; }
12
pow(double,int)13 __device__ double pow(double, int) { return 2.0; }
14
15 // HOST-DAG: call {{.*}}double @_Z3powdd(double noundef 1.000000e+00, double noundef 1.000000e+00)
16 double X = pow(1.0, 1);
17
cpow(double,double)18 constexpr double cpow(double, double) { return 11.0; }
19
cpow(double,int)20 constexpr __device__ double cpow(double, int) { return 12.0; }
21
22 // HOST-DAG: @CX = global double 1.100000e+01
23 double CX = cpow(11.0, 1);
24
25 // DEV-DAG: @CY = addrspace(1) externally_initialized global double 1.200000e+01
26 __device__ double CY = cpow(12.0, 1);
27
28 struct A {
powA29 double pow(double, double) { return 3.0; }
30
powA31 __device__ double pow(double, int) { return 4.0; }
32 };
33
34 A a;
35
36 // HOST-DAG: call {{.*}}double @_ZN1A3powEdd(ptr {{.*}}@a, double noundef 3.000000e+00, double noundef 1.000000e+00)
37 double AX = a.pow(3.0, 1);
38
39 struct CA {
cpowCA40 constexpr double cpow(double, double) const { return 13.0; }
41
cpowCA42 constexpr __device__ double cpow(double, int) const { return 14.0; }
43 };
44
45 const CA ca;
46
47 // HOST-DAG: @CAX = global double 1.300000e+01
48 double CAX = ca.cpow(13.0, 1);
49
50 // DEV-DAG: @CAY = addrspace(1) externally_initialized global double 1.400000e+01
51 __device__ double CAY = ca.cpow(14.0, 1);
52