https://github.com/abhinavgaba updated
https://github.com/llvm/llvm-project/pull/145454
>From d867ab3a20b1999cb5b00554fdb21cbe6766a1d5 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.g...@intel.com>
Date: Sun, 22 Jun 2025 18:36:16 -0700
Subject: [PATCH 1/3] [Clang][OpenMP] Capture mapped pointers on `target` by
reference.
For the following:
```c
int *p;
\#pragma omp target map(p[0]) // (A)
(void)p;
\#pragma omp target map(p) // (B)
(void)p;
\#pragma omp target map(p, p[0]) // (C)
(void)p;
\#pragma omp target map(p[0], p) // (D)
(void)p;
```
For (A), the pointer `p` is predetermined `firstprivate`, so it
should be (and is) captured by-copy. However, for (B), (C), and
(D), since `p` is already listed in a `map` clause, it's not
predetermined `firstprivate`, and hence, should be captured
by-reference, like any other mapped variable.
To ensure the correct handling of (C) and (D), the following
changes were made:
1. In SemaOpenMP, we now ensure that `p` is marked to be
captured by-reference in these cases.
2. We no longer ignore `map(p)` during codegen of `target`
constructs, even if there's another map like `map(p[0])` that
would have been mapped using a PTR_AND_OBJ map.
3. For cases like (D), we now handle `map(p)` before `map(p[0])`,
so the former gets the TARGET_PARAM flag and sets the kernel
argument.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 37 ++-
clang/lib/Sema/SemaOpenMP.cpp | 51 ++-
...arget_map_both_pointer_pointee_codegen.cpp | 290 +++++++++---------
.../mapping/map_ptr_and_subscript_global.c | 65 ++++
4 files changed, 287 insertions(+), 156 deletions(-)
create mode 100644 offload/test/mapping/map_ptr_and_subscript_global.c
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 8ccc37ef98a74..39571105e26b2 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7270,8 +7270,14 @@ class MappableExprsHandler {
// of arguments, hence MEMBER_OF(4)
//
// map(p, p[:100])
+ // For "pragma omp target":
+ // &p, &p, sizeof(p), TARGET_PARAM | TO | FROM
+ // &p, &p[0], 100*sizeof(float), PTR_AND_OBJ | TO | FROM (*)
+ // Otherwise:
// ===> map(p[:100])
// &p, &p[0], 100*sizeof(float), TARGET_PARAM | PTR_AND_OBJ | TO | FROM
+ // (*) We need to use PTR_AND_OBJ here to ensure that the mapped copies of
+ // p and p[0] get attached.
// Track if the map information being generated is the first for a capture.
bool IsCaptureFirstInfo = IsFirstComponentList;
@@ -7289,14 +7295,26 @@ class MappableExprsHandler {
// components.
bool IsExpressionFirstInfo = true;
bool FirstPointerInComplexData = false;
+ bool SkipStandalonePtrMapping = false;
Address BP = Address::invalid();
const Expr *AssocExpr = I->getAssociatedExpression();
const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
const auto *OASE = dyn_cast<ArraySectionExpr>(AssocExpr);
const auto *OAShE = dyn_cast<OMPArrayShapingExpr>(AssocExpr);
- if (AreBothBasePtrAndPteeMapped && std::next(I) == CE)
+ // For map(p, p[0]) on a "target" construct, we need to map "p" by itself
+ // as it has to be passed by-reference as the kernel argument.
+ // For other constructs, we can skip mapping "p" because the PTR_AND_OBJ
+ // mapping for map(p[0]) will take care of mapping p as well.
+ SkipStandalonePtrMapping =
+ AreBothBasePtrAndPteeMapped &&
+ (!isa<const OMPExecutableDirective *>(CurDir) ||
+ !isOpenMPTargetExecutionDirective(
+ cast<const OMPExecutableDirective
*>(CurDir)->getDirectiveKind()));
+
+ if (SkipStandalonePtrMapping && std::next(I) == CE)
return;
+
if (isa<MemberExpr>(AssocExpr)) {
// The base is the 'this' pointer. The content of the pointer is going
// to be the base of the field being mapped.
@@ -7672,7 +7690,7 @@ class MappableExprsHandler {
getMapTypeBits(MapType, MapModifiers, MotionModifiers,
IsImplicit,
!IsExpressionFirstInfo || RequiresReference ||
FirstPointerInComplexData ||
IsMemberReference,
- AreBothBasePtrAndPteeMapped ||
+ SkipStandalonePtrMapping ||
(IsCaptureFirstInfo && !RequiresReference),
IsNonContiguous);
@@ -8811,8 +8829,19 @@ class MappableExprsHandler {
++EI;
}
}
- llvm::stable_sort(DeclComponentLists, [](const MapData &LHS,
- const MapData &RHS) {
+ llvm::stable_sort(DeclComponentLists, [VD](const MapData &LHS,
+ const MapData &RHS) {
+ // For cases like map(p, p[0], p[0][0]), the shortest map, like map(p)
+ // in this case, should be handled first, to ensure that it gets the
+ // TARGET_PARAM flag.
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
+ std::get<0>(LHS);
+ OMPClauseMappableExprCommon::MappableExprComponentListRef ComponentsR =
+ std::get<0>(RHS);
+ if (VD && VD->getType()->isAnyPointerType() && Components.size() == 1 &&
+ ComponentsR.size() > 1)
+ return true;
+
ArrayRef<OpenMPMapModifierKind> MapModifiers = std::get<2>(LHS);
OpenMPMapClauseKind MapType = std::get<1>(RHS);
bool HasPresent =
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 00f4658180807..02e4e7b910d2e 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2146,6 +2146,7 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl
*D, unsigned Level,
// | ptr | n.a. | - | x | - | - |
bycopy|
// | ptr | n.a. | x | - | - | - | null
|
// | ptr | n.a. | - | - | - | x | byref
|
+ // | ptr | n.a. | - | - | - | x, x[] |
bycopy|
// | ptr | n.a. | - | - | - | x[] |
bycopy|
// | ptr | n.a. | - | - | x | |
bycopy|
// | ptr | n.a. | - | - | x | x |
bycopy|
@@ -2171,18 +2172,22 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl
*D, unsigned Level,
// - For pointers mapped by value that have either an implicit map or an
// array section, the runtime library may pass the NULL value to the
// device instead of the value passed to it by the compiler.
+ // - If both a pointer an a dereference of it are mapped, then the pointer
+ // should be passed by reference.
if (Ty->isReferenceType())
Ty = Ty->castAs<ReferenceType>()->getPointeeType();
- // Locate map clauses and see if the variable being captured is referred to
- // in any of those clauses. Here we only care about variables, not fields,
- // because fields are part of aggregates.
+ // Locate map clauses and see if the variable being captured is mapped by
+ // itself, or referred to, in any of those clauses. Here we only care about
+ // variables, not fields, because fields are part of aggregates.
bool IsVariableAssociatedWithSection = false;
+ bool IsVariableItselfMapped = false;
DSAStack->checkMappableExprComponentListsForDeclAtLevel(
D, Level,
[&IsVariableUsedInMapClause, &IsVariableAssociatedWithSection,
+ &IsVariableItselfMapped,
D](OMPClauseMappableExprCommon::MappableExprComponentListRef
MapExprComponents,
OpenMPClauseKind WhereFoundClauseKind) {
@@ -2198,8 +2203,19 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl
*D, unsigned Level,
assert(EI != EE && "Invalid map expression!");
- if (isa<DeclRefExpr>(EI->getAssociatedExpression()))
- IsVariableUsedInMapClause |= EI->getAssociatedDeclaration() == D;
+ if (isa<DeclRefExpr>(EI->getAssociatedExpression()) &&
+ EI->getAssociatedDeclaration() == D) {
+ IsVariableUsedInMapClause = true;
+
+ // If the component list has only one element, it's for mapping the
+ // variable itself, like map(p). This takes precedence in
+ // determining how it's captured, so we don't need to look further
+ // for any other maps that use the variable (like map(p[0]) etc.)
+ if (MapExprComponents.size() == 1) {
+ IsVariableItselfMapped = true;
+ return true;
+ }
+ }
++EI;
if (EI == EE)
@@ -2213,8 +2229,10 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl
*D, unsigned Level,
isa<MemberExpr>(EI->getAssociatedExpression()) ||
isa<OMPArrayShapingExpr>(Last->getAssociatedExpression())) {
IsVariableAssociatedWithSection = true;
- // There is nothing more we need to know about this variable.
- return true;
+ // We've found a case like map(p[0]) or map(p->a) or map(*p),
+ // so we are done with this particular map, but we need to keep
+ // looking in case we find a map(p).
+ return false;
}
// Keep looking for more map info.
@@ -2223,8 +2241,23 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl
*D, unsigned Level,
if (IsVariableUsedInMapClause) {
// If variable is identified in a map clause it is always captured by
- // reference except if it is a pointer that is dereferenced somehow.
- IsByRef = !(Ty->isPointerType() && IsVariableAssociatedWithSection);
+ // reference except if it is a pointer that is dereferenced somehow, but
+ // not itself mapped.
+ //
+ // OpenMP 6.0, 7.1.1: Data sharing attribute rules, variables referenced
+ // in a construct::
+ // If a list item in a has_device_addr clause or in a map clause on the
+ // target construct has a base pointer, and the base pointer is a scalar
+ // variable *that is not a list item in a map clause on the construct*,
+ // the base pointer is firstprivate.
+ //
+ // OpenMP 4.5, 2.15.1.1: Data-sharing Attribute Rules for Variables
+ // Referenced in a Construct:
+ // If an array section is a list item in a map clause on the target
+ // construct and the array section is derived from a variable for which
+ // the type is pointer then that variable is firstprivate.
+ IsByRef = IsVariableItselfMapped ||
+ !(Ty->isPointerType() && IsVariableAssociatedWithSection);
} else {
// By default, all the data that has a scalar type is mapped by copy
// (except for reduction variables).
diff --git a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
index 87fa7fe462daa..9a8f234da718c 100644
--- a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
+++ b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
@@ -1,174 +1,178 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex
"__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]"
"pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
// 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
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++
-std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple
i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - |
FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
+// CHECK: @.[[KERNEL00:__omp_offloading_.*foov_l[0-9]+]].region_id = weak
constant i8 0
+// CHECK: [[SIZE00:@.+]] = private unnamed_addr constant [2 x i64] [i64
{{8|4}}, i64 8]
+// CHECK: [[MYTYPE00:@.+]] = private unnamed_addr constant [2 x i64] [i64 35,
i64 19]
+
+// CHECK: @.[[KERNEL01:__omp_offloading_.*foov_l[0-9]+]].region_id = weak
constant i8 0
+// CHECK: [[SIZE01:@.+]] = private unnamed_addr constant [2 x i64] [i64
{{8|4}}, i64 4]
+// CHECK: [[MYTYPE01:@.+]] = private unnamed_addr constant [2 x i64] [i64 35,
i64 19]
+
+// CHECK: @.[[KERNEL02:__omp_offloading_.*foov_l[0-9]+]].region_id = weak
constant i8 0
+// CHECK: [[SIZE02:@.+]] = private unnamed_addr constant [2 x i64] [i64
{{8|4}}, i64 4]
+// CHECK: [[MYTYPE02:@.+]] = private unnamed_addr constant [2 x i64] [i64 35,
i64 19]
+
+// CHECK: [[SIZE03:@.+]] = private unnamed_addr constant [1 x i64] [i64 4]
+// CHECK: [[MYTYPE03:@.+]] = private unnamed_addr constant [1 x i64] [i64 51]
+
extern void *malloc (int __size) throw () __attribute__ ((__malloc__));
+// CHECK-LABEL: define{{.*}}@_Z3foov{{.*}}(
void foo() {
int *ptr = (int *) malloc(3 * sizeof(int));
+// Region 00
+// &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM
+// &ptr, &ptr[0], 2 * sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ
+//
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32
0, ptr @.[[KERNEL00]].region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0,
i32 2
+// CHECK-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
+// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32
3
+// CHECK-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
+// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+//
+// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]]
+// CHECK-DAG: store ptr [[VAR0]], ptr [[P0]]
+//
+// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0,
i{{.+}} 1
+// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0,
i{{.+}} 1
+// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP1]]
+// CHECK-DAG: store ptr [[RVAR00:%.+]], ptr [[P1]]
+//
+// CHECK-DAG: [[RVAR00]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}} 0
+// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]]
+//
+// CHECK-DAG: call void @[[KERNEL00]](ptr [[VAR0]])
#pragma omp target map(ptr, ptr[0:2])
{
ptr[1] = 6;
}
+
+// Region 01
+// &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM
+// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ
+//
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32
0, ptr @.[[KERNEL01]].region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0,
i32 2
+// CHECK-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
+// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32
3
+// CHECK-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
+// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+//
+// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]]
+// CHECK-DAG: store ptr [[VAR0]], ptr [[P0]]
+//
+// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0,
i{{.+}} 1
+// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0,
i{{.+}} 1
+// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP1]]
+// CHECK-DAG: store ptr [[RVAR02:%.+]], ptr [[P1]]
+//
+// CHECK-DAG: [[RVAR02]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}}
2
+// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]]
+//
+// CHECK-DAG: call void @[[KERNEL01]](ptr [[VAR0]])
#pragma omp target map(ptr, ptr[2])
{
ptr[2] = 8;
}
- #pragma omp target data map(ptr, ptr[2])
+
+// Region 02
+// &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM
+// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ
+//
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32
0, ptr @.[[KERNEL02]].region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0,
i32 2
+// CHECK-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
+// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32
3
+// CHECK-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
+// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+//
+// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]]
+// CHECK-DAG: store ptr [[VAR0]], ptr [[P0]]
+//
+// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0,
i{{.+}} 1
+// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0,
i{{.+}} 1
+// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP1]]
+// CHECK-DAG: store ptr [[RVAR02:%.+]], ptr [[P1]]
+//
+// CHECK-DAG: [[RVAR02]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}}
2
+// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]]
+//
+// CHECK-DAG: call void @[[KERNEL02]](ptr [[VAR0]])
+ #pragma omp target map(ptr[2], ptr)
{
ptr[2] = 9;
}
+
+// Region 03
+// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
+// FIXME: PARAM seems to be redundant here.
+//
+// CHECK-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1,
i32 1, ptr [[BPGEP:.+]], ptr [[PGEP:.+]], ptr [[SIZE03]], ptr [[MYTYPE03]], ptr
null, ptr null)
+// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+//
+// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0,
i{{.+}} 0
+// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]]
+// CHECK-DAG: store ptr [[RVAR02:%.+]], ptr [[P0]]
+//
+// CHECK-DAG: [[RVAR02]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}}
2
+// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]]
+ #pragma omp target data map(ptr, ptr[2])
+ {
+ ptr[2] = 10;
+ }
}
-#endif
-// CHECK-LABEL: define {{[^@]+}}@_Z3foov
-// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
+
+// CHECK-LABEL: define internal void
+// CHECK-SAME: @[[KERNEL00]](ptr {{[^,]*}}[[PTR:%[^,]+]])
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[PTR:%.*]] = alloca ptr, align 8
-// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca
[[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[KERNEL_ARGS5:%.*]] = alloca
[[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[CALL:%.*]] = call noalias noundef ptr @_Z6malloci(i32
noundef signext 12) #[[ATTR3:[0-9]+]]
-// CHECK-NEXT: store ptr [[CALL]], ptr [[PTR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr
[[TMP1]], i64 0
-// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP2]], align 8
-// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_PTRS]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
-// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
-// CHECK-NEXT: store ptr null, ptr [[TMP4]], align 8
-// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
-// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_PTRS]], i32 0, i32 0
-// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
-// CHECK-NEXT: store i32 3, ptr [[TMP7]], align 4
-// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
-// CHECK-NEXT: store i32 1, ptr [[TMP8]], align 4
-// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
-// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8
-// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
-// CHECK-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8
-// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
-// CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8
-// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
-// CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 8
-// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
-// CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8
-// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
-// CHECK-NEXT: store ptr null, ptr [[TMP14]], align 8
-// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
-// CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8
-// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
-// CHECK-NEXT: store i64 0, ptr [[TMP16]], align 8
-// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
-// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP17]], align
4
-// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
-// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4
-// CHECK-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
-// CHECK-NEXT: store i32 0, ptr [[TMP19]], align 4
-// CHECK-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr
@[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr
@.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15.region_id, ptr
[[KERNEL_ARGS]])
-// CHECK-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0
-// CHECK-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label
[[OMP_OFFLOAD_CONT:%.*]]
-// CHECK: omp_offload.failed:
-// CHECK-NEXT: call void
@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15(ptr [[TMP0]]) #[[ATTR3]]
-// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]
-// CHECK: omp_offload.cont:
-// CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[TMP23:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr
[[TMP23]], i64 2
-// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP24]], align 8
-// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_PTRS3]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[ARRAYIDX1]], ptr [[TMP25]], align 8
-// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_MAPPERS4]], i64 0, i64 0
-// CHECK-NEXT: store ptr null, ptr [[TMP26]], align 8
-// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0
-// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_PTRS3]], i32 0, i32 0
-// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0
-// CHECK-NEXT: store i32 3, ptr [[TMP29]], align 4
-// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1
-// CHECK-NEXT: store i32 1, ptr [[TMP30]], align 4
-// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2
-// CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8
-// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3
-// CHECK-NEXT: store ptr [[TMP28]], ptr [[TMP32]], align 8
-// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4
-// CHECK-NEXT: store ptr @.offload_sizes.1, ptr [[TMP33]], align 8
-// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5
-// CHECK-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP34]], align 8
-// CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6
-// CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8
-// CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7
-// CHECK-NEXT: store ptr null, ptr [[TMP36]], align 8
-// CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8
-// CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8
-// CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9
-// CHECK-NEXT: store i64 0, ptr [[TMP38]], align 8
-// CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10
-// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP39]], align
4
-// CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11
-// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP40]], align 4
-// CHECK-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12
-// CHECK-NEXT: store i32 0, ptr [[TMP41]], align 4
-// CHECK-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr
@[[GLOB1]], i64 -1, i32 -1, i32 0, ptr
@.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19.region_id, ptr
[[KERNEL_ARGS5]])
-// CHECK-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0
-// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED6:%.*]], label
[[OMP_OFFLOAD_CONT7:%.*]]
-// CHECK: omp_offload.failed6:
-// CHECK-NEXT: call void
@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19(ptr [[TMP22]]) #[[ATTR3]]
-// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT7]]
-// CHECK: omp_offload.cont7:
-// CHECK-NEXT: [[TMP44:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds i32, ptr
[[TMP44]], i64 2
-// CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP45]], align 8
-// CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_PTRS10]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[ARRAYIDX8]], ptr [[TMP46]], align 8
-// CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_MAPPERS11]], i64 0, i64 0
-// CHECK-NEXT: store ptr null, ptr [[TMP47]], align 8
-// CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0
-// CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_PTRS10]], i32 0, i32 0
-// CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1]],
i64 -1, i32 1, ptr [[TMP48]], ptr [[TMP49]], ptr @.offload_sizes.3, ptr
@.offload_maptypes.4, ptr null, ptr null)
-// CHECK-NEXT: [[TMP50:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds i32, ptr
[[TMP50]], i64 2
-// CHECK-NEXT: store i32 9, ptr [[ARRAYIDX12]], align 4
-// CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0
-// CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds [1 x ptr], ptr
[[DOTOFFLOAD_PTRS10]], i32 0, i32 0
-// CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64
-1, i32 1, ptr [[TMP51]], ptr [[TMP52]], ptr @.offload_sizes.3, ptr
@.offload_maptypes.4, ptr null, ptr null)
+// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr
+// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]]
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]]
+// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]]
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]],
i{{.*}} 1
+// CHECK-NEXT: store i32 6, ptr [[ARRAYIDX]]
// CHECK-NEXT: ret void
-//
-//
-// CHECK-LABEL: define
{{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15
-// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR2:[0-9]+]] {
+
+// CHECK-LABEL: define internal void
+// CHECK-SAME: @[[KERNEL01]](ptr {{[^,]*}}[[PTR:%[^,]+]])
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8
-// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]],
i64 1
-// CHECK-NEXT: store i32 6, ptr [[ARRAYIDX]], align 4
+// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr
+// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]]
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]]
+// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]]
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]],
i{{.*}} 2
+// CHECK-NEXT: store i32 8, ptr [[ARRAYIDX]]
// CHECK-NEXT: ret void
-//
-//
-// CHECK-LABEL: define
{{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19
-// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR2]] {
+
+// CHECK-LABEL: define internal void
+// CHECK-SAME: @[[KERNEL02]](ptr {{[^,]*}}[[PTR:%[^,]+]])
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8
-// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]],
i64 2
-// CHECK-NEXT: store i32 8, ptr [[ARRAYIDX]], align 4
+// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr
+// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]]
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]]
+// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]]
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]],
i{{.*}} 2
+// CHECK-NEXT: store i32 9, ptr [[ARRAYIDX]]
// CHECK-NEXT: ret void
-//
+#endif
diff --git a/offload/test/mapping/map_ptr_and_subscript_global.c
b/offload/test/mapping/map_ptr_and_subscript_global.c
new file mode 100644
index 0000000000000..1861064bc95df
--- /dev/null
+++ b/offload/test/mapping/map_ptr_and_subscript_global.c
@@ -0,0 +1,65 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+int x[10];
+int *p;
+
+void f1() {
+ p = &x[0];
+ p[0] = 111;
+ p[1] = 222;
+ p[2] = 333;
+ p[3] = 444;
+
+#pragma omp target enter data map(to : p)
+#pragma omp target enter data map(to:p[0 : 5])
+
+ int **p_mappedptr = (int **)omp_get_mapped_ptr(&p, omp_get_default_device());
+ int *p0_mappedptr =
+ (int *)omp_get_mapped_ptr(&p[0], omp_get_default_device());
+
+ printf("p_mappedptr %s null\n", p_mappedptr == (int **)NULL ? "==" : "!=");
+ printf("p0_mappedptr %s null\n", p0_mappedptr == (int *)NULL ? "==" : "!=");
+
+// CHECK: p_mappedptr != null
+// CHECK: p0_mappedptr != null
+
+// p is predetermined firstprivate, so its address will be different from
+// the mapped address for this construct. So, any changes to p within the
+// region will not be visible after the construct.
+#pragma omp target map(p[0]) firstprivate(p_mappedptr, p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[0]);
+ // CHECK: 111 0 1
+ p++;
+ }
+
+// For the remaining constructs, p is not firstprivate, so its address will
+// be the same as the mapped address, and changes to p will be visible to any
+// subsequent regions.
+#pragma omp target map(to : p[0], p) firstprivate(p_mappedptr, p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[0]);
+ // CHECK: 111 1 1
+ p++;
+ }
+
+#pragma omp target map(to : p, p[0]) firstprivate(p_mappedptr, p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[-1]);
+ // CHECK: 222 1 1
+ p++;
+ }
+
+#pragma omp target map(present, alloc : p) firstprivate(p_mappedptr,
p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[-2]);
+ // CHECK: 333 1 1
+ }
+
+#pragma omp target exit data map(delete:p[0 : 5], p)
+}
+
+int main() { f1(); }
>From 97b2dfff5b85e6e011eb241615fb4babd9bfb7ab Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.g...@intel.com>
Date: Mon, 23 Jun 2025 21:19:21 -0700
Subject: [PATCH 2/3] Add another test, minor clang-format fix.
---
.../mapping/map_ptr_and_subscript_global.c | 7 +-
.../mapping/map_ptr_and_subscript_local.c | 66 +++++++++++++++++++
2 files changed, 70 insertions(+), 3 deletions(-)
create mode 100644 offload/test/mapping/map_ptr_and_subscript_local.c
diff --git a/offload/test/mapping/map_ptr_and_subscript_global.c
b/offload/test/mapping/map_ptr_and_subscript_global.c
index 1861064bc95df..880a1e405b383 100644
--- a/offload/test/mapping/map_ptr_and_subscript_global.c
+++ b/offload/test/mapping/map_ptr_and_subscript_global.c
@@ -14,7 +14,7 @@ void f1() {
p[3] = 444;
#pragma omp target enter data map(to : p)
-#pragma omp target enter data map(to:p[0 : 5])
+#pragma omp target enter data map(to : p[0 : 5])
int **p_mappedptr = (int **)omp_get_mapped_ptr(&p, omp_get_default_device());
int *p0_mappedptr =
@@ -53,13 +53,14 @@ void f1() {
p++;
}
-#pragma omp target map(present, alloc : p) firstprivate(p_mappedptr,
p0_mappedptr)
+#pragma omp target map(present, alloc : p)
\
+ firstprivate(p_mappedptr, p0_mappedptr)
{
printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[-2]);
// CHECK: 333 1 1
}
-#pragma omp target exit data map(delete:p[0 : 5], p)
+#pragma omp target exit data map(delete : p[0 : 5], p)
}
int main() { f1(); }
diff --git a/offload/test/mapping/map_ptr_and_subscript_local.c
b/offload/test/mapping/map_ptr_and_subscript_local.c
new file mode 100644
index 0000000000000..24c5d0c7c5532
--- /dev/null
+++ b/offload/test/mapping/map_ptr_and_subscript_local.c
@@ -0,0 +1,66 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+int x[10];
+
+void f1() {
+ int *p;
+ p = &x[0];
+ p[0] = 111;
+ p[1] = 222;
+ p[2] = 333;
+ p[3] = 444;
+
+#pragma omp target enter data map(to : p)
+#pragma omp target enter data map(to : p[0 : 5])
+
+ int **p_mappedptr = (int **)omp_get_mapped_ptr(&p, omp_get_default_device());
+ int *p0_mappedptr =
+ (int *)omp_get_mapped_ptr(&p[0], omp_get_default_device());
+
+ printf("p_mappedptr %s null\n", p_mappedptr == (int **)NULL ? "==" : "!=");
+ printf("p0_mappedptr %s null\n", p0_mappedptr == (int *)NULL ? "==" : "!=");
+
+// CHECK: p_mappedptr != null
+// CHECK: p0_mappedptr != null
+
+// p is predetermined firstprivate, so its address will be different from
+// the mapped address for this construct. So, any changes to p within the
+// region will not be visible after the construct.
+#pragma omp target map(p[0]) firstprivate(p_mappedptr, p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[0]);
+ // CHECK: 111 0 1
+ p++;
+ }
+
+// For the remaining constructs, p is not firstprivate, so its address will
+// be the same as the mapped address, and changes to p will be visible to any
+// subsequent regions.
+#pragma omp target map(to : p[0], p) firstprivate(p_mappedptr, p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[0]);
+ // CHECK: 111 1 1
+ p++;
+ }
+
+#pragma omp target map(to : p, p[0]) firstprivate(p_mappedptr, p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[-1]);
+ // CHECK: 222 1 1
+ p++;
+ }
+
+#pragma omp target map(present, alloc : p)
\
+ firstprivate(p_mappedptr, p0_mappedptr)
+ {
+ printf("%d %d %d\n", p[0], p_mappedptr == &p, p0_mappedptr == &p[-2]);
+ // CHECK: 333 1 1
+ }
+
+#pragma omp target exit data map(delete : p[0 : 5], p)
+}
+
+int main() { f1(); }
>From a31ea2a50e0e3206a98bf1c3580343f609db4d63 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.g...@intel.com>
Date: Tue, 24 Jun 2025 11:21:01 -0700
Subject: [PATCH 3/3] Update the comparator for sort, plus minor additions to
offload tests.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 6 +++---
offload/test/mapping/map_ptr_and_subscript_global.c | 7 +++++++
offload/test/mapping/map_ptr_and_subscript_local.c | 7 +++++++
3 files changed, 17 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 39571105e26b2..f126feb214b6a 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8838,9 +8838,9 @@ class MappableExprsHandler {
std::get<0>(LHS);
OMPClauseMappableExprCommon::MappableExprComponentListRef ComponentsR =
std::get<0>(RHS);
- if (VD && VD->getType()->isAnyPointerType() && Components.size() == 1 &&
- ComponentsR.size() > 1)
- return true;
+ if (VD && VD->getType()->isAnyPointerType() &&
+ Components.size() != ComponentsR.size())
+ return Components.size() < ComponentsR.size();
ArrayRef<OpenMPMapModifierKind> MapModifiers = std::get<2>(LHS);
OpenMPMapClauseKind MapType = std::get<1>(RHS);
diff --git a/offload/test/mapping/map_ptr_and_subscript_global.c
b/offload/test/mapping/map_ptr_and_subscript_global.c
index 880a1e405b383..fdbd050be4cc7 100644
--- a/offload/test/mapping/map_ptr_and_subscript_global.c
+++ b/offload/test/mapping/map_ptr_and_subscript_global.c
@@ -60,6 +60,13 @@ void f1() {
// CHECK: 333 1 1
}
+ // The following map(from:p) should not bring back p, because p is an
+ // attached pointer. So, it should still point to the same original
+ // location, &x[0], on host.
+#pragma omp target exit data map(always, from : p)
+ printf("%d %d\n", p[0], p == &x[0]);
+ // CHECK: 111 1
+
#pragma omp target exit data map(delete : p[0 : 5], p)
}
diff --git a/offload/test/mapping/map_ptr_and_subscript_local.c
b/offload/test/mapping/map_ptr_and_subscript_local.c
index 24c5d0c7c5532..30ef0ce706be9 100644
--- a/offload/test/mapping/map_ptr_and_subscript_local.c
+++ b/offload/test/mapping/map_ptr_and_subscript_local.c
@@ -60,6 +60,13 @@ void f1() {
// CHECK: 333 1 1
}
+ // The following map(from:p) should not bring back p, because p is an
+ // attached pointer. So, it should still point to the same original
+ // location, &x[0], on host.
+#pragma omp target exit data map(always, from : p)
+ printf("%d %d\n", p[0], p == &x[0]);
+ // CHECK: 111 1
+
#pragma omp target exit data map(delete : p[0 : 5], p)
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
