1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
|
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
/// ========================================================================
/// Test: Firstprivate pointer handling in OpenMP target regions
/// ========================================================================
///
/// This test verifies that pointers with firstprivate semantics get the
/// OMP_MAP_LITERAL flag, enabling the runtime to pass pointer values directly
/// without performing present table lookups.
///
/// Map type values:
/// 288 = OMP_MAP_TARGET_PARAM (32) + OMP_MAP_LITERAL (256)
/// Used for explicit firstprivate(ptr)
///
/// 800 = OMP_MAP_TARGET_PARAM (32) + OMP_MAP_LITERAL (256) + OMP_MAP_IS_PTR (512)
/// Used for implicit firstprivate pointers (e.g., from defaultmap clauses)
/// Note: 512 is OMP_MAP_IS_PTR, not IMPLICIT. Implicitness is tracked separately.
///
/// 544 = OMP_MAP_TARGET_PARAM (32) + OMP_MAP_IS_PTR (512)
/// Incorrect behavior - missing LITERAL flag, causes runtime present table lookup
///
///==========================================================================
/// Test 1: Explicit firstprivate(pointer) → map type 288
///==========================================================================
// CHECK-DAG: @.offload_maptypes{{[^.]*}} = private unnamed_addr constant [1 x i64] [i64 288]
// CHECK-DAG: @.offload_sizes{{[^.]*}} = private unnamed_addr constant [1 x i64] zeroinitializer
void test1_explicit_firstprivate() {
double *ptr = nullptr;
// Explicit firstprivate should generate map type 288
// (TARGET_PARAM | LITERAL, no IS_PTR flag for explicit clauses)
#pragma omp target firstprivate(ptr)
{
if (ptr) ptr[0] = 1.0;
}
}
///==========================================================================
/// Test 2: defaultmap(firstprivate:pointer) → map type 800
///==========================================================================
// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 800]
// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
void test2_defaultmap_firstprivate_pointer() {
double *ptr = nullptr;
// defaultmap(firstprivate:pointer) creates implicit firstprivate
// Should generate map type 800 (TARGET_PARAM | LITERAL | IS_PTR)
#pragma omp target defaultmap(firstprivate:pointer)
{
if (ptr) ptr[0] = 2.0;
}
}
///==========================================================================
/// Test 3: defaultmap(firstprivate:scalar) with double → map type 800
///==========================================================================
// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 800]
void test3_defaultmap_scalar_double() {
double d = 3.0;
// OpenMP's "scalar" category excludes pointers but includes arithmetic types
// Double gets implicit firstprivate → map type 800
#pragma omp target defaultmap(firstprivate:scalar)
{
d += 1.0;
}
}
///==========================================================================
/// Test 4: Pointer with defaultmap(firstprivate:scalar) → map type 800
///==========================================================================
// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 800]
// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
void test4_pointer_with_scalar_defaultmap() {
double *ptr = nullptr;
// Note: defaultmap(firstprivate:scalar) does NOT apply to pointers (scalar excludes pointers).
// However, the pointer still gets 800 because in OpenMP 5.0+, pointers without explicit
// data-sharing attributes are implicitly firstprivate and lowered as IS_PTR|LITERAL|TARGET_PARAM.
// This is the default pointer behavior, NOT due to the scalar defaultmap.
#pragma omp target defaultmap(firstprivate:scalar)
{
if (ptr) ptr[0] = 4.0;
}
}
///==========================================================================
/// Test 5: Multiple pointers with explicit firstprivate → all get 288
///==========================================================================
// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 288]
// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [3 x i64] zeroinitializer
void test5_multiple_firstprivate() {
int *a = nullptr;
float *b = nullptr;
double *c = nullptr;
// All explicit firstprivate pointers get map type 288
#pragma omp target firstprivate(a, b, c)
{
if (a) a[0] = 6;
if (b) b[0] = 7.0f;
if (c) c[0] = 8.0;
}
}
///==========================================================================
/// Test 6: Pointer to const with firstprivate → map type 288
///==========================================================================
// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 288]
// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
void test6_const_pointer() {
const double *const_ptr = nullptr;
// Const pointer with explicit firstprivate → 288
#pragma omp target firstprivate(const_ptr)
{
if (const_ptr) {
double val = const_ptr[0];
(void)val;
}
}
}
///==========================================================================
/// Test 7: Pointer-to-pointer with firstprivate → map type 288
///==========================================================================
// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 288]
// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
void test7_pointer_to_pointer() {
int **pp = nullptr;
// Pointer-to-pointer with explicit firstprivate → 288
#pragma omp target firstprivate(pp)
{
if (pp && *pp) (*pp)[0] = 9;
}
}
///==========================================================================
/// Verification: The key fix is that firstprivate pointers now include
/// the LITERAL flag (256), which tells the runtime to pass the pointer
/// value directly instead of performing a present table lookup.
///
/// Before fix: Pointers got 544 (TARGET_PARAM | IS_PTR) → runtime lookup
/// After fix: Pointers get 288 or 800 (includes LITERAL) → direct pass
///==========================================================================
#endif // HEADER
|
