[Clang][OpenMP] Capture mapped pointers on `target` by reference. by abhinavgaba · Pull Request #145454 · llvm/llvm-project (original) (raw)

@llvm/pr-subscribers-offload
@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-clang

Author: Abhinav Gaba (abhinavgaba)

Changes

For the following:

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:

In SemaOpenMP, we now ensure that p is marked to be captured by-reference in these cases.
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.
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 --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], 100sizeof(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(AssocExpr); const auto *OASE = dyn_cast(AssocExpr); const auto *OAShE = dyn_cast(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()->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(EI->getAssociatedExpression()) || isa(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]

[Clang][OpenMP] Capture mapped pointers on target by reference. by abhinavgaba · Pull Request #145454 · llvm/llvm-project (original) (raw)

[Clang][OpenMP] Capture mapped pointers on `target` by reference. by abhinavgaba · Pull Request #145454 · llvm/llvm-project (original) (raw)