[clang] [llvm] [Clang][OpenMP] Capture mapped pointers on `target` by reference. (PR #145454)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Jun 24 08:32:44 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-clang
Author: Abhinav Gaba (abhinavgaba)
<details>
<summary>Changes</summary>
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.
---
Patch is 34.47 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/145454.diff
5 Files Affected:
- (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+33-4)
- (modified) clang/lib/Sema/SemaOpenMP.cpp (+42-9)
- (modified) clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp (+147-143)
- (added) offload/test/mapping/map_ptr_and_subscript_global.c (+66)
- (added) offload/test/mapping/map_ptr_and_subscript_local.c (+66)
``````````diff
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:%.*]] = getelement...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/145454
More information about the cfe-commits
mailing list