Skip to content

Commit 73f219f

Browse files
Jason-VanBeusekomJason-VanBeusekom
committed
Update lit tests for variable and non weak aliasing
1 parent 16c1a68 commit 73f219f

File tree

4 files changed

+128
-54
lines changed

4 files changed

+128
-54
lines changed

‎clang/lib/CodeGen/CodeGenModule.cpp‎

Lines changed: 6 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -4054,17 +4054,12 @@ static bool shouldSkipAliasEmission(const CodeGenModule &CGM,
40544054

40554055
// Check if the aliasee exists.
40564056
if (!CGM.lookupRepresentativeDecl(AA->getAliasee(), AliaseeGD)) {
4057-
if (LangOpts.CUDA)
4058-
// In CUDA/HIP, if the aliasee is not found, skip the alias emission.
4059-
// This is not a hard error as this branch is executed for both the host
4060-
// and device, with no respect to where the aliasee is defined.
4061-
return true;
4062-
4063-
// For OpenMP, lookupRepresentativeDecl should always find the aliasee, this
4064-
// is an error
4065-
CGM.getDiags().Report(AA->getLocation(), diag::err_alias_to_undefined)
4066-
<< false << true;
4067-
return false;
4057+
// If the aliasee is not found, skip the alias emission.
4058+
// This is not a hard error as this branch is executed for both the host
4059+
// and device, with no respect to where the aliasee is defined.
4060+
// For some OpenMP cases (functions) this will return true even if the
4061+
// aliasee is not on the device, which is handled by the case below
4062+
return true;
40684063
}
40694064

40704065
const auto *AliaseeDecl = dyn_cast<ValueDecl>(AliaseeGD.getDecl());

‎clang/test/CodeGenHIP/hip_weak_alias.cpp‎

Lines changed: 53 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -8,40 +8,54 @@
88
#define __host__ __attribute__((host))
99

1010
extern "C" {
11-
1211
//.
12+
// HOST: @__HostVar = global i32 1, align 4
1313
// HOST: @__hip_cuid_ = global i8 0
1414
// HOST: @llvm.compiler.used = appending global [1 x ptr] [ptr @__hip_cuid_], section "llvm.metadata"
1515
// HOST: @HostFunc = weak alias i32 (), ptr @__HostFunc
16+
// HOST: @HostFunc_ = alias i32 (), ptr @__HostFunc
17+
// HOST: @HostVar = weak alias i32, ptr @__HostVar
18+
// HOST: @HostVar_ = alias i32, ptr @__HostVar
1619
// HOST: @Two = weak alias i32 (), ptr @__Two
17-
// HOST: @Four = weak alias i32 (), ptr @__Four
20+
// HOST: @Two_ = alias i32 (), ptr @__Two
21+
// HOST: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
22+
// HOST: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
23+
// HOST: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
24+
// HOST: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
1825
//.
1926
// DEVICE: @__hip_cuid_ = addrspace(1) global i8 0
2027
// DEVICE: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata"
2128
// DEVICE: @One = weak alias i32 (), ptr @__One
29+
// DEVICE: @One_ = alias i32 (), ptr @__One
2230
// DEVICE: @Two = weak alias i32 (), ptr @__Two
23-
// DEVICE: @Three = weak alias i32 (), ptr @__Three
24-
// DEVICE: @Five = weak alias i32 (), ptr @__Five
25-
// DEVICE: @_Z3Sixv = weak alias i32 (), ptr @_Z5__Sixv
26-
// DEVICE: @_Z3Sixf = weak alias float (float), ptr @_Z5__Sixf
31+
// DEVICE: @Two_ = alias i32 (), ptr @__Two
32+
// DEVICE: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
33+
// DEVICE: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
34+
// DEVICE: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
35+
// DEVICE: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
2736
//.
2837
// HOST-LABEL: define dso_local i32 @__HostFunc(
2938
// HOST-SAME: ) #[[ATTR0:[0-9]+]] {
3039
// HOST-NEXT: [[ENTRY:.*:]]
3140
// HOST-NEXT: ret i32 42
3241
//
3342
int __HostFunc(void) { return 42; }
43+
int __HostVar = 1;
3444
int HostFunc(void) __attribute__((weak, alias("__HostFunc")));
45+
int HostFunc_(void) __attribute__((alias("__HostFunc")));
46+
extern int __attribute__((weak, alias("__HostVar"))) HostVar;
47+
extern int __attribute__((alias("__HostVar"))) HostVar_;
3548

3649
// DEVICE-LABEL: define dso_local i32 @__One(
3750
// DEVICE-SAME: ) #[[ATTR0:[0-9]+]] {
3851
// DEVICE-NEXT: [[ENTRY:.*:]]
3952
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
4053
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
41-
// DEVICE-NEXT: ret i32 2
54+
// DEVICE-NEXT: ret i32 1
4255
//
43-
__device__ int __One(void) { return 2; }
56+
__device__ int __One(void) { return 1; }
4457
__device__ int One(void) __attribute__((weak, alias("__One")));
58+
__device__ int One_(void) __attribute__((alias("__One")));
4559

4660
// HOST-LABEL: define dso_local i32 @__Two(
4761
// HOST-SAME: ) #[[ATTR0]] {
@@ -57,45 +71,49 @@ __device__ int One(void) __attribute__((weak, alias("__One")));
5771
//
5872
__host__ __device__ int __Two(void) { return 2; }
5973
__host__ __device__ int Two(void) __attribute__((weak, alias("__Two")));
74+
__host__ __device__ int Two_(void) __attribute__((alias("__Two")));
75+
}
6076

61-
// DEVICE-LABEL: define linkonce_odr i32 @__Three(
62-
// DEVICE-SAME: ) #[[ATTR0]] comdat {
63-
// DEVICE-NEXT: [[ENTRY:.*:]]
64-
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
65-
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
66-
// DEVICE-NEXT: ret i32 2
67-
//
68-
__device__ constexpr int __Three(void) { return 2; }
69-
__device__ int Three(void) __attribute__((weak, alias("__Three")));
70-
71-
// HOST-LABEL: define linkonce_odr i32 @__Four(
77+
// HOST-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
7278
// HOST-SAME: ) #[[ATTR0]] comdat {
7379
// HOST-NEXT: [[ENTRY:.*:]]
74-
// HOST-NEXT: ret i32 2
80+
// HOST-NEXT: ret i32 5
7581
//
76-
constexpr int __Four(void) { return 2; }
77-
int Four(void) __attribute__((weak, alias("__Four")));
78-
79-
// DEVICE-LABEL: define linkonce_odr i32 @__Five(
82+
// DEVICE-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
8083
// DEVICE-SAME: ) #[[ATTR0]] comdat {
8184
// DEVICE-NEXT: [[ENTRY:.*:]]
8285
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
8386
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
84-
// DEVICE-NEXT: ret i32 2
87+
// DEVICE-NEXT: ret i32 5
8588
//
86-
__device__ constexpr int __Five(void) { return 2; }
87-
__device__ int Five(void) __attribute__((weak, alias("__Five")));
88-
}
89+
__host__ __device__ constexpr int __Three(void) { return 5; }
90+
__host__ __device__ int Three(void) __attribute__((weak, alias("_Z7__Threev")));
91+
__host__ __device__ int Three_(void) __attribute__((alias("_Z7__Threev")));
8992

90-
// DEVICE-LABEL: define dso_local noundef i32 @_Z5__Sixv(
93+
94+
// HOST-LABEL: define dso_local noundef i32 @_Z6__Fourv(
95+
// HOST-SAME: ) #[[ATTR0]] {
96+
// HOST-NEXT: [[ENTRY:.*:]]
97+
// HOST-NEXT: ret i32 2
98+
//
99+
// DEVICE-LABEL: define dso_local noundef i32 @_Z6__Fourv(
91100
// DEVICE-SAME: ) #[[ATTR0]] {
92101
// DEVICE-NEXT: [[ENTRY:.*:]]
93102
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
94103
// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
95104
// DEVICE-NEXT: ret i32 2
96105
//
97-
__device__ int __Six(void) { return 2; }
98-
// DEVICE-LABEL: define dso_local noundef float @_Z5__Sixf(
106+
__host__ __device__ int __Four(void) { return 2; }
107+
// HOST-LABEL: define dso_local noundef float @_Z6__Fourf(
108+
// HOST-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
109+
// HOST-NEXT: [[ENTRY:.*:]]
110+
// HOST-NEXT: [[F_ADDR:%.*]] = alloca float, align 4
111+
// HOST-NEXT: store float [[F]], ptr [[F_ADDR]], align 4
112+
// HOST-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4
113+
// HOST-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
114+
// HOST-NEXT: ret float [[MUL]]
115+
//
116+
// DEVICE-LABEL: define dso_local noundef float @_Z6__Fourf(
99117
// DEVICE-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
100118
// DEVICE-NEXT: [[ENTRY:.*:]]
101119
// DEVICE-NEXT: [[RETVAL:%.*]] = alloca float, align 4, addrspace(5)
@@ -107,9 +125,10 @@ __device__ int __Six(void) { return 2; }
107125
// DEVICE-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
108126
// DEVICE-NEXT: ret float [[MUL]]
109127
//
110-
__device__ float __Six(float f) { return 2.0f * f; }
111-
__device__ int Six(void) __attribute__((weak, alias("_Z5__Sixv")));
112-
__device__ float Six(float f) __attribute__((weak, alias("_Z5__Sixf")));
128+
__host__ __device__ float __Four(float f) { return 2.0f * f; }
129+
__host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));
130+
__host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));
131+
113132
//.
114133
// HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
115134
//.

‎clang/test/OpenMP/amdgcn_weak_alias.c‎

Lines changed: 53 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -6,18 +6,43 @@
66
// RUN: %clang_cc1 -fopenmp -x c -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=DEVICE
77

88
//.
9+
// HOST: @__One_var = global i32 1, align 4
10+
// HOST: @__Two_var = global i32 2, align 4
11+
// HOST: @__Three_var = global i32 3, align 4
12+
// HOST: @.offloading.entry_name = internal unnamed_addr constant [10 x i8] c"__Two_var\00", section ".llvm.rodata.offloading", align 1
13+
// HOST: @.offloading.entry.__Two_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Two_var, ptr @.offloading.entry_name, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8
14+
// HOST: @.offloading.entry_name.1 = internal unnamed_addr constant [12 x i8] c"__Three_var\00", section ".llvm.rodata.offloading", align 1
15+
// HOST: @.offloading.entry.__Three_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Three_var, ptr @.offloading.entry_name.1, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8
916
// HOST: @One = weak alias i32 (), ptr @__One
17+
// HOST: @One_ = alias i32 (), ptr @__One
18+
// HOST: @One_var = weak alias i32, ptr @__One_var
19+
// HOST: @One_var_ = alias i32, ptr @__One_var
1020
// HOST: @Two = weak alias i32 (), ptr @__Two
21+
// HOST: @Two_ = alias i32 (), ptr @__Two
22+
// HOST: @Two_var = weak alias i32, ptr @__Two_var
23+
// HOST: @Two_var_ = alias i32, ptr @__Two_var
1124
// HOST: @Three = weak alias i32 (), ptr @__Three
25+
// HOST: @Three_ = alias i32 (), ptr @__Three
26+
// HOST: @Three_var = weak alias i32, ptr @__Three_var
27+
// HOST: @Three_var_ = alias i32, ptr @__Three_var
1228
//.
1329
// DEVICE: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0
1430
// DEVICE: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0
1531
// DEVICE: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden addrspace(1) constant i32 0
1632
// DEVICE: @__omp_rtl_assume_no_thread_state = weak_odr hidden addrspace(1) constant i32 0
1733
// DEVICE: @__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0
34+
// DEVICE: @__Two_var = addrspace(1) global i32 2, align 4
35+
// DEVICE: @__Three_var = addrspace(1) global i32 3, align 4
1836
// DEVICE: @Two = weak hidden alias i32 (), ptr @__Two
37+
// DEVICE: @Two_ = hidden alias i32 (), ptr @__Two
38+
// DEVICE: @Two_var = weak alias i32, addrspacecast (ptr addrspace(1) @__Two_var to ptr)
39+
// DEVICE: @Two_var_ = alias i32, addrspacecast (ptr addrspace(1) @__Two_var to ptr)
1940
// DEVICE: @Three = weak hidden alias i32 (), ptr @__Three
2041
// DEVICE: @Three.1 = weak hidden alias i32 (), ptr @__Three
42+
// DEVICE: @Three_ = hidden alias i32 (), ptr @__Three
43+
// DEVICE: @Three_.2 = hidden alias i32 (), ptr @__Three
44+
// DEVICE: @Three_var = weak alias i32, addrspacecast (ptr addrspace(1) @__Three_var to ptr)
45+
// DEVICE: @Three_var_ = alias i32, addrspacecast (ptr addrspace(1) @__Three_var to ptr)
2146
//.
2247
// HOST-LABEL: define dso_local i32 @__One(
2348
// HOST-SAME: ) #[[ATTR0:[0-9]+]] {
@@ -26,6 +51,11 @@
2651
//
2752
int __One(void) { return 1; }
2853
int One(void) __attribute__ ((weak, alias("__One")));
54+
int One_(void) __attribute__ ((alias("__One")));
55+
56+
int __One_var = 1;
57+
extern int __attribute__((weak, alias("__One_var"))) One_var;
58+
extern int __attribute__((alias("__One_var"))) One_var_;
2959

3060
#pragma omp declare target
3161
// HOST-LABEL: define dso_local i32 @__Two(
@@ -42,6 +72,11 @@ int One(void) __attribute__ ((weak, alias("__One")));
4272
//
4373
int __Two(void) { return 2; }
4474
int Two(void) __attribute__ ((weak, alias("__Two")));
75+
int Two_(void) __attribute__ ((alias("__Two")));
76+
77+
int __Two_var = 2;
78+
extern int __attribute__((weak, alias("__Two_var"))) Two_var;
79+
extern int __attribute__((alias("__Two_var"))) Two_var_;
4580
#pragma omp end declare target
4681

4782
#pragma omp declare target
@@ -58,20 +93,30 @@ int Two(void) __attribute__ ((weak, alias("__Two")));
5893
// DEVICE-NEXT: ret i32 3
5994
//
6095
int __Three(void) { return 3; }
96+
int __Three_var = 3;
6197
#pragma omp end declare target
6298
int Three(void) __attribute__ ((weak, alias("__Three")));
99+
int Three_(void) __attribute__ ((alias("__Three")));
100+
extern int __attribute__((weak, alias("__Three_var"))) Three_var;
101+
extern int __attribute__((alias("__Three_var"))) Three_var_;
63102
//.
64103
// HOST: attributes #[[ATTR0]] = { noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
65104
//.
66105
// DEVICE: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
67106
//.
68-
// HOST: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
69-
// HOST: [[META1:![0-9]+]] = !{i32 7, !"openmp", i32 51}
70-
// HOST: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
107+
// HOST: [[META0:![0-9]+]] = !{i32 1, !"__Two_var", i32 0, i32 0}
108+
// HOST: [[META1:![0-9]+]] = !{i32 1, !"__Three_var", i32 0, i32 1}
109+
// HOST: [[META2:![0-9]+]] = !{ptr @.offloading.entry_name}
110+
// HOST: [[META3:![0-9]+]] = !{ptr @.offloading.entry_name.1}
111+
// HOST: [[META4:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
112+
// HOST: [[META5:![0-9]+]] = !{i32 7, !"openmp", i32 51}
113+
// HOST: [[META6:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
71114
//.
72-
// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
73-
// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
74-
// DEVICE: [[META2:![0-9]+]] = !{i32 7, !"openmp", i32 51}
75-
// DEVICE: [[META3:![0-9]+]] = !{i32 7, !"openmp-device", i32 51}
76-
// DEVICE: [[META4:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
115+
// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"__Two_var", i32 0, i32 0}
116+
// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"__Three_var", i32 0, i32 1}
117+
// DEVICE: [[META2:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
118+
// DEVICE: [[META3:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
119+
// DEVICE: [[META4:![0-9]+]] = !{i32 7, !"openmp", i32 51}
120+
// DEVICE: [[META5:![0-9]+]] = !{i32 7, !"openmp-device", i32 51}
121+
// DEVICE: [[META6:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
77122
//.

‎clang/test/OpenMP/amdgcn_weak_alias.cpp‎

Lines changed: 16 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,10 +6,16 @@
66
//.
77
// HOST: @_Z3Onev = weak alias i32 (), ptr @_Z5__Onev
88
// HOST: @_Z3Onef = weak alias float (float), ptr @_Z5__Onef
9+
// HOST: @_Z4One_v = alias i32 (), ptr @_Z5__Onev
10+
// HOST: @_Z4One_f = alias float (float), ptr @_Z5__Onef
911
// HOST: @_Z3Twov = weak alias i32 (), ptr @_Z5__Twov
1012
// HOST: @_Z3Twof = weak alias float (float), ptr @_Z5__Twof
13+
// HOST: @_Z4Two_v = alias i32 (), ptr @_Z5__Twov
14+
// HOST: @_Z4Two_f = alias float (float), ptr @_Z5__Twof
1115
// HOST: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
16+
// HOST: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
1217
// HOST: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
18+
// HOST: @_Z5Four_v = alias i32 (), ptr @_Z6__Fourv
1319
//.
1420
// DEVICE: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0
1521
// DEVICE: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0
@@ -18,15 +24,17 @@
1824
// DEVICE: @__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0
1925
// DEVICE: @_Z3Twov = weak hidden alias i32 (), ptr @_Z5__Twov
2026
// DEVICE: @_Z3Twof = weak hidden alias float (float), ptr @_Z5__Twof
27+
// DEVICE: @_Z4Two_v = hidden alias i32 (), ptr @_Z5__Twov
28+
// DEVICE: @_Z4Two_f = hidden alias float (float), ptr @_Z5__Twof
2129
// DEVICE: @_Z5Threev = weak hidden alias i32 (), ptr @_Z7__Threev
30+
// DEVICE: @_Z6Three_v = hidden alias i32 (), ptr @_Z7__Threev
2231
//.
2332
// HOST-LABEL: define dso_local noundef i32 @_Z5__Onev(
2433
// HOST-SAME: ) #[[ATTR0:[0-9]+]] {
2534
// HOST-NEXT: [[ENTRY:.*:]]
2635
// HOST-NEXT: ret i32 1
2736
//
2837
int __One(void) { return 1; }
29-
3038
// HOST-LABEL: define dso_local noundef float @_Z5__Onef(
3139
// HOST-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
3240
// HOST-NEXT: [[ENTRY:.*:]]
@@ -39,6 +47,8 @@ int __One(void) { return 1; }
3947
float __One(float f) { return 1.0f * f; }
4048
int One(void) __attribute__((weak, alias("_Z5__Onev")));
4149
float One(float f) __attribute__((weak, alias("_Z5__Onef")));
50+
int One_(void) __attribute__((alias("_Z5__Onev")));
51+
float One_(float f) __attribute__((alias("_Z5__Onef")));
4252

4353
#pragma omp declare target
4454
// HOST-LABEL: define dso_local noundef i32 @_Z5__Twov(
@@ -78,6 +88,8 @@ int __Two(void) { return 2; }
7888
float __Two(float f) { return 2.0f * f; }
7989
int Two(void) __attribute__((weak, alias("_Z5__Twov")));
8090
float Two(float f) __attribute__((weak, alias("_Z5__Twof")));
91+
int Two_(void) __attribute__((alias("_Z5__Twov")));
92+
float Two_(float f) __attribute__((alias("_Z5__Twof")));
8193
#pragma omp end declare target
8294

8395
#pragma omp declare target
@@ -95,14 +107,17 @@ float Two(float f) __attribute__((weak, alias("_Z5__Twof")));
95107
//
96108
constexpr int __Three(void) { return 3; }
97109
int Three(void) __attribute__((weak, alias("_Z7__Threev")));
110+
int Three_(void) __attribute__((alias("_Z7__Threev")));
98111
#pragma omp end declare target
112+
99113
// HOST-LABEL: define linkonce_odr noundef i32 @_Z6__Fourv(
100114
// HOST-SAME: ) #[[ATTR0]] comdat {
101115
// HOST-NEXT: [[ENTRY:.*:]]
102116
// HOST-NEXT: ret i32 4
103117
//
104118
constexpr int __Four(void) { return 4; }
105119
int Four(void) __attribute__((weak, alias("_Z6__Fourv")));
120+
int Four_(void) __attribute__((alias("_Z6__Fourv")));
106121
//.
107122
// HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
108123
//.

0 commit comments

Comments
 (0)