diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index a5f2f0efa2c3b..28818d5ac0568 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7275,8 +7275,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; @@ -7294,14 +7300,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(CurDir) || + !isOpenMPTargetExecutionDirective( + cast(CurDir)->getDirectiveKind())); + + if (SkipStandalonePtrMapping && std::next(I) == CE) return; + if (isa(AssocExpr)) { // The base is the 'this' pointer. The content of the pointer is going // to be the base of the field being mapped. @@ -7677,7 +7695,7 @@ class MappableExprsHandler { getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit, !IsExpressionFirstInfo || RequiresReference || FirstPointerInComplexData || IsMemberReference, - AreBothBasePtrAndPteeMapped || + SkipStandalonePtrMapping || (IsCaptureFirstInfo && !RequiresReference), IsNonContiguous); @@ -8863,8 +8881,56 @@ class MappableExprsHandler { CurCaptureVarInfo.append(CurInfoForComponentLists); }; - GenerateInfoForComponentLists(DeclComponentLists, + // Next, we break-down the lists of components into lists that should be + // handled together. + // + // For now, we handle maps on pointers by themselves, and everything else + // together. + // + // TODO: Do this based on which lists have the same attachable-base-pointer + // e.g. The map clauses below, which are present on the same construct, + // should be handled grouped together based on their + // attachable-base-pointers: + // map-clause | attachable-base-pointer + // --------------------------+------------------------ + // map(p, ps) | none + // map(p[0]) | p + // map(p[0]->b, p[0]->c) | p[0] + // map(ps->d, ps->e, ps->pt) | ps + // map(ps->pt->d, ps->pt->e) | ps->pt + + MapDataArrayTy ListsThatMapPointerVDItself; + MapDataArrayTy RemainingLists; + bool IsVDPointerType = VD && VD->getType()->isPointerType(); + + for (const MapData &L : DeclComponentLists) { + if (IsVDPointerType) { + OMPClauseMappableExprCommon::MappableExprComponentListRef Components = + std::get<0>(L); + bool IsMapOfPointerVD = + Components.size() == 1 && + Components[0].getAssociatedDeclaration() && + Components[0].getAssociatedDeclaration()->getCanonicalDecl() == VD; + if (IsMapOfPointerVD) { + ListsThatMapPointerVDItself.push_back(L); + continue; + } + } + RemainingLists.push_back(L); + } + + // If VD is a pointer, and there are component-lists mapping VD itself, + // like: `int *p; ... map(p)`, we handle them first, and + // the first one from the should get the TARGET_PARAM flag. + GenerateInfoForComponentLists(ListsThatMapPointerVDItself, /*IsEligibleForTargetParamFlag=*/true); + // Then we handle all the other lists together. Note that since we already + // added TARGET_PARAM for the pointer case, we shouldn't add it again if + // if there are other lists that get handled here, for example, the list + // for `map(p[0:10]` in `int *p; ... map(p, p[0:10])`. + GenerateInfoForComponentLists( + RemainingLists, + /*IsEligibleForTargetParamFlag=*/ListsThatMapPointerVDItself.empty()); } /// Generate the base pointers, section pointers, sizes, map types, and diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index a30acbe9a4bca..db07bd61bf406 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(EI->getAssociatedExpression())) - IsVariableUsedInMapClause |= EI->getAssociatedDeclaration() == D; + if (isa(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 4264e58dd2152..8a167a19ddbc9 100644 --- a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp +++ b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp @@ -23,22 +23,25 @@ void f2() { void f3() { int *ptr; - // &ptr, &ptr[0], sizeof(ptr[0:2]), TO | FROM | PARAM | PTR_AND_OBJ + // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr, &ptr[0], 2 * sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ #pragma omp target map(ptr, ptr[0:2]) ptr[1] = 7; } void f4() { int *ptr; - // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ + // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ #pragma omp target map(ptr, ptr[2]) ptr[2] = 8; } void f5() { int *ptr; - // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ - #pragma omp target map(ptr[2], ptr) + // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ + #pragma omp target map(ptr[2], ptr) ptr[2] = 9; } @@ -63,12 +66,12 @@ void f7() { // CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]] // CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4] // CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]] -// CHECK: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 8] -// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] -// CHECK: @.offload_sizes.5 = private unnamed_addr constant [1 x i64] [i64 4] -// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] -// CHECK: @.offload_sizes.7 = private unnamed_addr constant [1 x i64] [i64 4] -// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] +// CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 8] +// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] +// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] +// CHECK: @.offload_sizes.7 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] // CHECK: @.offload_sizes.9 = private unnamed_addr constant [1 x i64] [i64 4] // CHECK: @.offload_maptypes.10 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] // CHECK: @.offload_sizes.11 = private unnamed_addr constant [1 x i64] [i64 4] @@ -131,25 +134,31 @@ void f7() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 0 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP0]], i64 0 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PTR]], ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr [[PTR]], ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[PTR]], ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // -// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l27 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l28 +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 1 // CHECK: store i32 7, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // @@ -158,25 +167,31 @@ void f7() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PTR]], ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr [[PTR]], ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[PTR]], ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // -// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l34 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l36 +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 // CHECK: store i32 8, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // @@ -185,25 +200,31 @@ void f7() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PTR]], ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr [[PTR]], ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[PTR]], ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // -// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l41 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l44 +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 // CHECK: store i32 9, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // diff --git a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen_global.cpp b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen_global.cpp index ee9cb9af4d3ca..e39d602dca316 100644 --- a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen_global.cpp +++ b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen_global.cpp @@ -22,20 +22,23 @@ void f2() { } void f3() { - // &ptr, &ptr[0], sizeof(ptr[0:2]), TO | FROM | PARAM | PTR_AND_OBJ + // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr, &ptr[0], 2 * sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ #pragma omp target map(ptr, ptr[0:2]) ptr[1] = 7; } void f4() { - // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ + // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ #pragma omp target map(ptr, ptr[2]) ptr[2] = 8; } void f5() { - // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ - #pragma omp target map(ptr[2], ptr) + // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ + #pragma omp target map(ptr[2], ptr) ptr[2] = 9; } @@ -58,12 +61,12 @@ void f7() { // CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]] // CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4] // CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] -// CHECK: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 8] -// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] -// CHECK: @.offload_sizes.5 = private unnamed_addr constant [1 x i64] [i64 4] -// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] -// CHECK: @.offload_sizes.7 = private unnamed_addr constant [1 x i64] [i64 4] -// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] +// CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 8] +// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] +// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] +// CHECK: @.offload_sizes.7 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] // CHECK: @.offload_sizes.9 = private unnamed_addr constant [1 x i64] [i64 4] // CHECK: @.offload_maptypes.10 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] // CHECK: @.offload_sizes.11 = private unnamed_addr constant [1 x i64] [i64 4] @@ -125,25 +128,31 @@ void f7() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 0 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP0]], i64 0 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ptr, ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr @ptr, ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr @ptr, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // -// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l26 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l27 +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 1 // CHECK: store i32 7, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // @@ -152,25 +161,31 @@ void f7() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ptr, ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr @ptr, ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr @ptr, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // -// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l32 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l34 +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 // CHECK: store i32 8, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // @@ -179,25 +194,31 @@ void f7() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ptr, ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr @ptr, ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr @ptr, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // -// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l38 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l41 +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 // CHECK: store i32 9, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // diff --git a/clang/test/OpenMP/target_map_ptr_and_star_global.cpp b/clang/test/OpenMP/target_map_ptr_and_star_global.cpp index 84899cb8e4fad..40d3cd24ac08f 100644 --- a/clang/test/OpenMP/target_map_ptr_and_star_global.cpp +++ b/clang/test/OpenMP/target_map_ptr_and_star_global.cpp @@ -28,8 +28,8 @@ void f3() { } void f4() { - // &ptr, &ptr[0], sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ | PARAM // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr, &ptr[0], sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ #pragma omp target map(*ptr, ptr) ptr[2] = 8; } @@ -42,8 +42,8 @@ void f4() { // CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]] // CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] // CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] -// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 4, i64 8] -// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x33]], i64 [[#0x3]]] +// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]] //. // CHECK-LABEL: define {{[^@]+}}@_Z2f1v // CHECK-SAME: () #[[ATTR0:[0-9]+]] { @@ -100,30 +100,30 @@ void f4() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ptr, ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr @ptr, ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr @ptr, ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr @ptr, ptr [[TMP5]], align 8 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[TMP1]], ptr [[TMP6]], align 8 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP7]], align 8 -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr @ptr, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l26 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META11]], !align [[META12]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 1 // CHECK: store i32 6, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // @@ -132,30 +132,30 @@ void f4() { // CHECK-SAME: () #[[ATTR0]] { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8 -// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ptr, ptr [[TMP1]], align 8 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr @ptr, ptr [[TMP2]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr @ptr, ptr [[TMP5]], align 8 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr @ptr, ptr [[TMP6]], align 8 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP7]], align 8 -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr @ptr, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l33 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META11]], !align [[META12]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 // CHECK: store i32 8, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // diff --git a/clang/test/OpenMP/target_map_ptr_and_star_local.cpp b/clang/test/OpenMP/target_map_ptr_and_star_local.cpp index 246c0c5f99a68..2eff1727e0e46 100644 --- a/clang/test/OpenMP/target_map_ptr_and_star_local.cpp +++ b/clang/test/OpenMP/target_map_ptr_and_star_local.cpp @@ -31,8 +31,8 @@ void f3() { void f4() { int *ptr; - // &ptr[0], &ptr[0], sizeof(ptr[0]), TO | FROM // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM + // &ptr[0], &ptr[0], sizeof(ptr[0]), TO | FROM #pragma omp target map(*ptr, ptr) ptr[2] = 8; } @@ -45,7 +45,7 @@ void f4() { // CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]] // CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] // CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x3]]] -// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 4, i64 8] +// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] // CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x3]]] //. // CHECK-LABEL: define {{[^@]+}}@_Z2f1v @@ -105,30 +105,30 @@ void f4() { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PTR]], ptr [[TMP2]], align 8 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 // CHECK: store ptr [[PTR]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[PTR]], ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP5]], align 8 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 // CHECK: store ptr [[TMP1]], ptr [[TMP6]], align 8 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[TMP2]], ptr [[TMP7]], align 8 -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP8]], align 8 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP7]], align 8 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l28 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META11]], !align [[META12]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 1 // CHECK: store i32 6, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // @@ -138,30 +138,30 @@ void f4() { // CHECK: entry: // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP3]], align 8 -// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP2]], ptr [[TMP4]], align 8 -// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP5]], align 8 -// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr [[PTR]], ptr [[TMP6]], align 8 -// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[PTR]], ptr [[TMP7]], align 8 -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP8]], align 8 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PTR]], ptr [[TMP2]], align 8 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PTR]], ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP1]], ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP7]], align 8 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 // // // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l36 -// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] { +// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1]] { // CHECK: entry: // CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8 -// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META11]], !align [[META12]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2 // CHECK: store i32 8, ptr [[ARRAYIDX]], align 4 // CHECK: ret void // diff --git a/clang/test/OpenMP/target_map_structptr_and_member_global.cpp b/clang/test/OpenMP/target_map_structptr_and_member_global.cpp index 523f88dc8dba3..dc2df3849f299 100644 --- a/clang/test/OpenMP/target_map_structptr_and_member_global.cpp +++ b/clang/test/OpenMP/target_map_structptr_and_member_global.cpp @@ -28,26 +28,24 @@ void f2() { } void f3() { - // &ps[0], &ps[0], sizeof(ps[0]), PARAM | ALLOC - // &ps, &ps, sizeof(ps), TO | MEMBER_OF(1) - // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(1) + // &ps, &ps, sizeof(ps), TO | PARAM + // &ps[0], &ps->y, sizeof(ps->y), TO #pragma omp target map(to: ps, ps->y) ps->y = 7; } void f4() { - // &ps[0], &ps[0], sizeof(ps[0]), PARAM | ALLOC - // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(1) - // &ps, &ps, sizeof(ps), TO | MEMBER_OF(1) + // &ps, &ps, sizeof(ps), TO | PARAM + // &ps[0], &ps->y, sizeof(ps->y), TO #pragma omp target map(to: ps->y, ps) ps->y = 8; } void f5() { - // &ps[0], &ps[0], sizeof(ps[0]), PARAM | ALLOC - // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(1) - // &ps, &ps, sizeof(ps), TO | MEMBER_OF(1) - // &ps[0], &ps->x, sizeof(ps->x), TO | MEMBER_OF(1) + // &ps, &ps, sizeof(ps), TO | PARAM + // &ps[0], &ps[0].x, ((&ps[0].y + 1) - &ps[0].x)/8, ALLOC + // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(2) + // &ps[0], &ps->x, sizeof(ps->x), TO | MEMBER_OF(2) #pragma omp target map(to: ps->y, ps, ps->x) ps->y = 9; } @@ -58,12 +56,12 @@ void f5() { // CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x21]]] // CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4] // CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x21]]] -// CHECK: @.offload_sizes.3 = private unnamed_addr constant [3 x i64] [i64 0, i64 8, i64 4] -// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]]] -// CHECK: @.offload_sizes.5 = private unnamed_addr constant [3 x i64] [i64 0, i64 4, i64 8] -// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]]] -// CHECK: @.offload_sizes.7 = private unnamed_addr constant [4 x i64] [i64 0, i64 4, i64 8, i64 2] -// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]]] +// CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x21]], i64 [[#0x1]]] +// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x21]], i64 [[#0x1]]] +// CHECK: @.offload_sizes.7 = private unnamed_addr constant [4 x i64] [i64 8, i64 0, i64 4, i64 2] +// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [4 x i64] [i64 [[#0x21]], i64 [[#0x0]], i64 [[#0x2000000000001]], i64 [[#0x2000000000001]]] //. // CHECK-LABEL: define dso_local void @_Z2f1v( // CHECK-SAME: ) #[[ATTR0:[0-9]+]] { @@ -123,46 +121,31 @@ void f5() { // CHECK: [[ENTRY:.*:]] // CHECK: [[TMP0:%.*]] = load ptr, ptr @ps, align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr @ps, align 8 -// CHECK: [[TMP2:%.*]] = load ptr, ptr @ps, align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP2]], i32 0, i32 1 -// CHECK: [[TMP3:%.*]] = getelementptr [[STRUCT_S]], ptr [[TMP1]], i32 1 -// CHECK: [[TMP4:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK: [[TMP5:%.*]] = ptrtoint ptr [[TMP1]] to i64 -// CHECK: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] -// CHECK: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) -// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES:%.*]], ptr align 8 @.offload_sizes.3, i64 24, i1 false) -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP8]], align 8 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP9]], align 8 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: store i64 [[TMP7]], ptr [[TMP10]], align 8 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP11]], align 8 -// CHECK: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr @ps, ptr [[TMP12]], align 8 -// CHECK: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr @ps, ptr [[TMP13]], align 8 -// CHECK: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP14]], align 8 -// CHECK: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 -// CHECK: store ptr [[TMP1]], ptr [[TMP15]], align 8 -// CHECK: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 -// CHECK: store ptr [[Y]], ptr [[TMP16]], align 8 -// CHECK: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 -// CHECK: store ptr null, ptr [[TMP17]], align 8 -// CHECK: [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP20:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 -// -// -// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l34( -// CHECK-SAME: ptr noundef [[PS:%.*]]) #[[ATTR1]] { +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ps, ptr [[TMP2]], align 8 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ps, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[Y]], ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP7]], align 8 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l33( +// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(8) [[PS:%.*]]) #[[ATTR1]] { // CHECK: [[ENTRY:.*:]] // CHECK: store ptr [[PS]], ptr [[PS_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: store i32 7, ptr [[Y]], align 4 // CHECK: ret void // @@ -172,46 +155,31 @@ void f5() { // CHECK: [[ENTRY:.*:]] // CHECK: [[TMP0:%.*]] = load ptr, ptr @ps, align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr @ps, align 8 -// CHECK: [[TMP2:%.*]] = load ptr, ptr @ps, align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP2]], i32 0, i32 1 -// CHECK: [[TMP3:%.*]] = getelementptr [[STRUCT_S]], ptr [[TMP1]], i32 1 -// CHECK: [[TMP4:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK: [[TMP5:%.*]] = ptrtoint ptr [[TMP1]] to i64 -// CHECK: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] -// CHECK: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) -// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES:%.*]], ptr align 8 @.offload_sizes.5, i64 24, i1 false) -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP8]], align 8 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP9]], align 8 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: store i64 [[TMP7]], ptr [[TMP10]], align 8 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP11]], align 8 -// CHECK: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr [[TMP1]], ptr [[TMP12]], align 8 -// CHECK: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[Y]], ptr [[TMP13]], align 8 -// CHECK: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP14]], align 8 -// CHECK: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 -// CHECK: store ptr @ps, ptr [[TMP15]], align 8 -// CHECK: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 -// CHECK: store ptr @ps, ptr [[TMP16]], align 8 -// CHECK: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 -// CHECK: store ptr null, ptr [[TMP17]], align 8 -// CHECK: [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP20:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 -// -// -// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l42( -// CHECK-SAME: ptr noundef [[PS:%.*]]) #[[ATTR1]] { +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ps, ptr [[TMP2]], align 8 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ps, ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[Y]], ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP7]], align 8 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l40( +// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(8) [[PS:%.*]]) #[[ATTR1]] { // CHECK: [[ENTRY:.*:]] // CHECK: store ptr [[PS]], ptr [[PS_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: store i32 8, ptr [[Y]], align 4 // CHECK: ret void // @@ -221,55 +189,55 @@ void f5() { // CHECK: [[ENTRY:.*:]] // CHECK: [[TMP0:%.*]] = load ptr, ptr @ps, align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr @ps, align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: [[TMP2:%.*]] = load ptr, ptr @ps, align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP2]], i32 0, i32 1 // CHECK: [[TMP3:%.*]] = load ptr, ptr @ps, align 8 -// CHECK: [[TMP4:%.*]] = load ptr, ptr @ps, align 8 -// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[TMP4]], i32 0, i32 0 -// CHECK: [[TMP5:%.*]] = getelementptr [[STRUCT_S]], ptr [[TMP1]], i32 1 -// CHECK: [[TMP6:%.*]] = ptrtoint ptr [[TMP5]] to i64 -// CHECK: [[TMP7:%.*]] = ptrtoint ptr [[TMP1]] to i64 -// CHECK: [[TMP8:%.*]] = sub i64 [[TMP6]], [[TMP7]] -// CHECK: [[TMP9:%.*]] = sdiv exact i64 [[TMP8]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) +// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[TMP3]], i32 0, i32 0 +// CHECK: [[TMP4:%.*]] = getelementptr i32, ptr [[Y]], i32 1 +// CHECK: [[TMP5:%.*]] = ptrtoint ptr [[TMP4]] to i64 +// CHECK: [[TMP6:%.*]] = ptrtoint ptr [[X]] to i64 +// CHECK: [[TMP7:%.*]] = sub i64 [[TMP5]], [[TMP6]] +// CHECK: [[TMP8:%.*]] = sdiv exact i64 [[TMP7]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) // CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES:%.*]], ptr align 8 @.offload_sizes.7, i64 32, i1 false) -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP10]], align 8 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP11]], align 8 -// CHECK: [[TMP12:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: store i64 [[TMP9]], ptr [[TMP12]], align 8 -// CHECK: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP13]], align 8 -// CHECK: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr [[TMP1]], ptr [[TMP14]], align 8 -// CHECK: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[Y]], ptr [[TMP15]], align 8 -// CHECK: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP16]], align 8 -// CHECK: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 -// CHECK: store ptr @ps, ptr [[TMP17]], align 8 -// CHECK: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 -// CHECK: store ptr @ps, ptr [[TMP18]], align 8 -// CHECK: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 -// CHECK: store ptr null, ptr [[TMP19]], align 8 -// CHECK: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 -// CHECK: store ptr [[TMP3]], ptr [[TMP20]], align 8 -// CHECK: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3 -// CHECK: store ptr [[X]], ptr [[TMP21]], align 8 -// CHECK: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3 -// CHECK: store ptr null, ptr [[TMP22]], align 8 -// CHECK: [[TMP23:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP24:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP25:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 -// -// -// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l51( -// CHECK-SAME: ptr noundef [[PS:%.*]]) #[[ATTR1]] { +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ps, ptr [[TMP9]], align 8 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr @ps, ptr [[TMP10]], align 8 +// CHECK: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP11]], align 8 +// CHECK: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP12]], align 8 +// CHECK: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[X]], ptr [[TMP13]], align 8 +// CHECK: [[TMP14:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK: store i64 [[TMP8]], ptr [[TMP14]], align 8 +// CHECK: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP15]], align 8 +// CHECK: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK: store ptr [[TMP0]], ptr [[TMP16]], align 8 +// CHECK: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK: store ptr [[Y]], ptr [[TMP17]], align 8 +// CHECK: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 +// CHECK: store ptr null, ptr [[TMP18]], align 8 +// CHECK: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 +// CHECK: store ptr [[TMP2]], ptr [[TMP19]], align 8 +// CHECK: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3 +// CHECK: store ptr [[X]], ptr [[TMP20]], align 8 +// CHECK: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3 +// CHECK: store ptr null, ptr [[TMP21]], align 8 +// CHECK: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP23:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP24:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l49( +// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(8) [[PS:%.*]]) #[[ATTR1]] { // CHECK: [[ENTRY:.*:]] // CHECK: store ptr [[PS]], ptr [[PS_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: store i32 9, ptr [[Y]], align 4 // CHECK: ret void // diff --git a/clang/test/OpenMP/target_map_structptr_and_member_local.cpp b/clang/test/OpenMP/target_map_structptr_and_member_local.cpp index b366f331941b7..44a982680bb8b 100644 --- a/clang/test/OpenMP/target_map_structptr_and_member_local.cpp +++ b/clang/test/OpenMP/target_map_structptr_and_member_local.cpp @@ -29,28 +29,26 @@ void f2() { void f3() { S s, *ps; - // &ps[0], &ps[0], sizeof(ps[0]), PARAM | ALLOC - // &ps, &ps, sizeof(ps), TO | MEMBER_OF(1) - // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(1) + // &ps, &ps, sizeof(ps), TO | PARAM + // &ps[0], &ps->y, sizeof(ps->y), TO #pragma omp target map(to: ps, ps->y) ps->y = 7; } void f4() { S s, *ps; - // &ps[0], &ps[0], sizeof(ps[0]), PARAM | ALLOC - // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(1) - // &ps, &ps, sizeof(ps), TO | MEMBER_OF(1) + // &ps, &ps, sizeof(ps), TO | PARAM + // &ps[0], &ps->y, sizeof(ps->y), TO #pragma omp target map(to: ps->y, ps) ps->y = 8; } void f5() { S s, *ps; - // &ps[0], &ps[0], sizeof(ps[0]), PARAM | ALLOC - // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(1) - // &ps, &ps, sizeof(ps), TO | MEMBER_OF(1) - // &ps[0], &ps->x, sizeof(ps->x), TO | MEMBER_OF(1) + // &ps, &ps, sizeof(ps), TO | PARAM + // &ps[0], &ps[0].x, ((&ps[0].y + 1) - &ps[0].x)/8, ALLOC + // &ps[0], &ps->y, sizeof(ps->y), TO | MEMBER_OF(2) + // &ps[0], &ps->x, sizeof(ps->x), TO | MEMBER_OF(2) #pragma omp target map(to: ps->y, ps, ps->x) ps->y = 9; } @@ -61,12 +59,12 @@ void f5() { // CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x21]]] // CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4] // CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x21]]] -// CHECK: @.offload_sizes.3 = private unnamed_addr constant [3 x i64] [i64 0, i64 8, i64 4] -// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]]] -// CHECK: @.offload_sizes.5 = private unnamed_addr constant [3 x i64] [i64 0, i64 4, i64 8] -// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]]] -// CHECK: @.offload_sizes.7 = private unnamed_addr constant [4 x i64] [i64 0, i64 4, i64 8, i64 2] -// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]]] +// CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x21]], i64 [[#0x1]]] +// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 8, i64 4] +// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x21]], i64 [[#0x1]]] +// CHECK: @.offload_sizes.7 = private unnamed_addr constant [4 x i64] [i64 8, i64 0, i64 4, i64 2] +// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [4 x i64] [i64 [[#0x21]], i64 [[#0x0]], i64 [[#0x2000000000001]], i64 [[#0x2000000000001]]] //. // CHECK-LABEL: define dso_local void @_Z2f1v( // CHECK-SAME: ) #[[ATTR0:[0-9]+]] { @@ -126,46 +124,31 @@ void f5() { // CHECK: [[ENTRY:.*:]] // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS:%.*]], align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr [[PS]], align 8 -// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PS]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP2]], i32 0, i32 1 -// CHECK: [[TMP3:%.*]] = getelementptr [[STRUCT_S]], ptr [[TMP1]], i32 1 -// CHECK: [[TMP4:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK: [[TMP5:%.*]] = ptrtoint ptr [[TMP1]] to i64 -// CHECK: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] -// CHECK: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) -// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES:%.*]], ptr align 8 @.offload_sizes.3, i64 24, i1 false) -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP8]], align 8 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP9]], align 8 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: store i64 [[TMP7]], ptr [[TMP10]], align 8 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP11]], align 8 -// CHECK: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr [[PS]], ptr [[TMP12]], align 8 -// CHECK: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[PS]], ptr [[TMP13]], align 8 -// CHECK: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP14]], align 8 -// CHECK: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 -// CHECK: store ptr [[TMP1]], ptr [[TMP15]], align 8 -// CHECK: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 -// CHECK: store ptr [[Y]], ptr [[TMP16]], align 8 -// CHECK: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 -// CHECK: store ptr null, ptr [[TMP17]], align 8 -// CHECK: [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP20:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 -// -// -// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l35( -// CHECK-SAME: ptr noundef [[PS:%.*]]) #[[ATTR1]] { +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PS]], ptr [[TMP2]], align 8 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PS]], ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[Y]], ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP7]], align 8 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l34( +// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(8) [[PS:%.*]]) #[[ATTR1]] { // CHECK: [[ENTRY:.*:]] // CHECK: store ptr [[PS]], ptr [[PS_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: store i32 7, ptr [[Y]], align 4 // CHECK: ret void // @@ -175,46 +158,31 @@ void f5() { // CHECK: [[ENTRY:.*:]] // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS:%.*]], align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr [[PS]], align 8 -// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PS]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP2]], i32 0, i32 1 -// CHECK: [[TMP3:%.*]] = getelementptr [[STRUCT_S]], ptr [[TMP1]], i32 1 -// CHECK: [[TMP4:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK: [[TMP5:%.*]] = ptrtoint ptr [[TMP1]] to i64 -// CHECK: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] -// CHECK: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) -// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES:%.*]], ptr align 8 @.offload_sizes.5, i64 24, i1 false) -// CHECK: [[TMP8:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP8]], align 8 -// CHECK: [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP9]], align 8 -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: store i64 [[TMP7]], ptr [[TMP10]], align 8 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP11]], align 8 -// CHECK: [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr [[TMP1]], ptr [[TMP12]], align 8 -// CHECK: [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[Y]], ptr [[TMP13]], align 8 -// CHECK: [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP14]], align 8 -// CHECK: [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 -// CHECK: store ptr [[PS]], ptr [[TMP15]], align 8 -// CHECK: [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 -// CHECK: store ptr [[PS]], ptr [[TMP16]], align 8 -// CHECK: [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 -// CHECK: store ptr null, ptr [[TMP17]], align 8 -// CHECK: [[TMP18:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP19:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP20:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: [[TMP21:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 -// -// -// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l44( -// CHECK-SAME: ptr noundef [[PS:%.*]]) #[[ATTR1]] { +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PS]], ptr [[TMP2]], align 8 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PS]], ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[Y]], ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP7]], align 8 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l42( +// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(8) [[PS:%.*]]) #[[ATTR1]] { // CHECK: [[ENTRY:.*:]] // CHECK: store ptr [[PS]], ptr [[PS_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: store i32 8, ptr [[Y]], align 4 // CHECK: ret void // @@ -224,55 +192,55 @@ void f5() { // CHECK: [[ENTRY:.*:]] // CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS:%.*]], align 8 // CHECK: [[TMP1:%.*]] = load ptr, ptr [[PS]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: [[TMP2:%.*]] = load ptr, ptr [[PS]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP2]], i32 0, i32 1 // CHECK: [[TMP3:%.*]] = load ptr, ptr [[PS]], align 8 -// CHECK: [[TMP4:%.*]] = load ptr, ptr [[PS]], align 8 -// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[TMP4]], i32 0, i32 0 -// CHECK: [[TMP5:%.*]] = getelementptr [[STRUCT_S]], ptr [[TMP1]], i32 1 -// CHECK: [[TMP6:%.*]] = ptrtoint ptr [[TMP5]] to i64 -// CHECK: [[TMP7:%.*]] = ptrtoint ptr [[TMP1]] to i64 -// CHECK: [[TMP8:%.*]] = sub i64 [[TMP6]], [[TMP7]] -// CHECK: [[TMP9:%.*]] = sdiv exact i64 [[TMP8]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) +// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[TMP3]], i32 0, i32 0 +// CHECK: [[TMP4:%.*]] = getelementptr i32, ptr [[Y]], i32 1 +// CHECK: [[TMP5:%.*]] = ptrtoint ptr [[TMP4]] to i64 +// CHECK: [[TMP6:%.*]] = ptrtoint ptr [[X]] to i64 +// CHECK: [[TMP7:%.*]] = sub i64 [[TMP5]], [[TMP6]] +// CHECK: [[TMP8:%.*]] = sdiv exact i64 [[TMP7]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) // CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES:%.*]], ptr align 8 @.offload_sizes.7, i64 32, i1 false) -// CHECK: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP10]], align 8 -// CHECK: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 -// CHECK: store ptr [[TMP1]], ptr [[TMP11]], align 8 -// CHECK: [[TMP12:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: store i64 [[TMP9]], ptr [[TMP12]], align 8 -// CHECK: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 -// CHECK: store ptr null, ptr [[TMP13]], align 8 -// CHECK: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 -// CHECK: store ptr [[TMP1]], ptr [[TMP14]], align 8 -// CHECK: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 -// CHECK: store ptr [[Y]], ptr [[TMP15]], align 8 -// CHECK: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 -// CHECK: store ptr null, ptr [[TMP16]], align 8 -// CHECK: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 -// CHECK: store ptr [[PS]], ptr [[TMP17]], align 8 -// CHECK: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 -// CHECK: store ptr [[PS]], ptr [[TMP18]], align 8 -// CHECK: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 -// CHECK: store ptr null, ptr [[TMP19]], align 8 -// CHECK: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 -// CHECK: store ptr [[TMP3]], ptr [[TMP20]], align 8 -// CHECK: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3 -// CHECK: store ptr [[X]], ptr [[TMP21]], align 8 -// CHECK: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3 -// CHECK: store ptr null, ptr [[TMP22]], align 8 -// CHECK: [[TMP23:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK: [[TMP24:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK: [[TMP25:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 -// CHECK: [[TMP26:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 -// -// -// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l54( -// CHECK-SAME: ptr noundef [[PS:%.*]]) #[[ATTR1]] { +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PS]], ptr [[TMP9]], align 8 +// CHECK: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[PS]], ptr [[TMP10]], align 8 +// CHECK: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr null, ptr [[TMP11]], align 8 +// CHECK: [[TMP12:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[TMP0]], ptr [[TMP12]], align 8 +// CHECK: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[X]], ptr [[TMP13]], align 8 +// CHECK: [[TMP14:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK: store i64 [[TMP8]], ptr [[TMP14]], align 8 +// CHECK: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP15]], align 8 +// CHECK: [[TMP16:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK: store ptr [[TMP0]], ptr [[TMP16]], align 8 +// CHECK: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK: store ptr [[Y]], ptr [[TMP17]], align 8 +// CHECK: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 +// CHECK: store ptr null, ptr [[TMP18]], align 8 +// CHECK: [[TMP19:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 +// CHECK: store ptr [[TMP2]], ptr [[TMP19]], align 8 +// CHECK: [[TMP20:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3 +// CHECK: store ptr [[X]], ptr [[TMP20]], align 8 +// CHECK: [[TMP21:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3 +// CHECK: store ptr null, ptr [[TMP21]], align 8 +// CHECK: [[TMP22:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP23:%.*]] = getelementptr inbounds [4 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: [[TMP24:%.*]] = getelementptr inbounds [4 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK: [[TMP25:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0 +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l52( +// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(8) [[PS:%.*]]) #[[ATTR1]] { // CHECK: [[ENTRY:.*:]] // CHECK: store ptr [[PS]], ptr [[PS_ADDR:%.*]], align 8 -// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8 -// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 1 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PS_ADDR]], align 8, !nonnull [[META13]], !align [[META14]] +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP1]], i32 0, i32 1 // CHECK: store i32 9, ptr [[Y]], align 4 // CHECK: ret void // diff --git a/offload/test/mapping/map_ptr_and_star_global.c b/offload/test/mapping/map_ptr_and_star_global.c index c3b0dd2f49e6b..255e2dee6512e 100644 --- a/offload/test/mapping/map_ptr_and_star_global.c +++ b/offload/test/mapping/map_ptr_and_star_global.c @@ -46,8 +46,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0], x0_hostaddr == &p[0]); - // EXPECTED: 111 1 1 0 - // CHECK: 111 0 1 0 + // CHECK: 111 1 1 0 p++; } @@ -56,8 +55,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-1], x0_hostaddr == &p[-1]); - // EXPECTED: 222 1 1 0 - // CHECK: {{[0-9]+}} 0 0 0 + // CHECK: 222 1 1 0 p++; } @@ -66,8 +64,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-2], x0_hostaddr == &p[-2]); - // EXPECTED: 333 1 1 0 - // CHECK: 111 1 0 0 + // CHECK: 333 1 1 0 } // The following map(from:p) should not bring back p, because p is an diff --git a/offload/test/mapping/map_ptr_and_star_local.c b/offload/test/mapping/map_ptr_and_star_local.c index f0ca84d1cc4dd..0d2d689ec9912 100644 --- a/offload/test/mapping/map_ptr_and_star_local.c +++ b/offload/test/mapping/map_ptr_and_star_local.c @@ -47,7 +47,7 @@ void f1() { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0], x0_hostaddr == &p[0]); // EXPECTED: 111 1 1 0 - // CHECK: 111 0 1 0 + // CHECK: 111 1 0 1 p++; } @@ -57,7 +57,7 @@ void f1() { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-1], x0_hostaddr == &p[-1]); // EXPECTED: 222 1 1 0 - // CHECK: {{[0-9]+}} 0 0 0 + // CHECK: 222 1 0 1 p++; } @@ -67,7 +67,7 @@ void f1() { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-2], x0_hostaddr == &p[-2]); // EXPECTED: 333 1 1 0 - // CHECK: 111 1 0 0 + // CHECK: 333 1 0 1 } // The following map(from:p) should not bring back p, because p is an @@ -75,7 +75,8 @@ void f1() { // 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 + // EXPECTED: 111 1 + // CHECK: 333 0 #pragma omp target exit data map(delete : p[0 : 5], p) } diff --git a/offload/test/mapping/map_ptr_and_subscript_global.c b/offload/test/mapping/map_ptr_and_subscript_global.c index a3a10b6c9b212..63486773a68ec 100644 --- a/offload/test/mapping/map_ptr_and_subscript_global.c +++ b/offload/test/mapping/map_ptr_and_subscript_global.c @@ -46,8 +46,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0], x0_hostaddr == &p[0]); - // EXPECTED: 111 1 1 0 - // CHECK: 111 0 1 0 + // CHECK: 111 1 1 0 p++; } @@ -56,8 +55,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-1], x0_hostaddr == &p[-1]); - // EXPECTED: 222 1 1 0 - // CHECK: 111 0 0 0 + // CHECK: 222 1 1 0 p++; } @@ -66,8 +64,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-2], x0_hostaddr == &p[-2]); - // EXPECTED: 333 1 1 0 - // CHECK: 111 1 0 0 + // CHECK: 333 1 1 0 } // The following map(from:p) should not bring back p, because p is an diff --git a/offload/test/mapping/map_ptr_and_subscript_local.c b/offload/test/mapping/map_ptr_and_subscript_local.c index bb44999541a7b..59c05fa876d96 100644 --- a/offload/test/mapping/map_ptr_and_subscript_local.c +++ b/offload/test/mapping/map_ptr_and_subscript_local.c @@ -46,8 +46,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0], x0_hostaddr == &p[0]); - // EXPECTED: 111 1 1 0 - // CHECK: 111 0 1 0 + // CHECK: 111 1 1 0 p++; } @@ -56,8 +55,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-1], x0_hostaddr == &p[-1]); - // EXPECTED: 222 1 1 0 - // CHECK: 111 0 0 0 + // CHECK: 222 1 1 0 p++; } @@ -66,8 +64,7 @@ void f1() { { printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-2], x0_hostaddr == &p[-2]); - // EXPECTED: 333 1 1 0 - // CHECK: 111 1 0 0 + // CHECK: 333 1 1 0 } // The following map(from:p) should not bring back p, because p is an diff --git a/offload/test/mapping/map_structptr_and_member_global.c b/offload/test/mapping/map_structptr_and_member_global.c index 10e72e070dbc5..f2fae005b79ab 100644 --- a/offload/test/mapping/map_structptr_and_member_global.c +++ b/offload/test/mapping/map_structptr_and_member_global.c @@ -52,7 +52,7 @@ void f1() { printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps->x, s0_hostaddr == &ps->x); // EXPECTED: 111 1 1 0 - // CHECK: 111 0 1 0 + // CHECK: 111 1 0 1 ps++; } @@ -62,7 +62,7 @@ void f1() { printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps[-1].x, s0_hostaddr == &ps[-1].x); // EXPECTED: 222 1 1 0 - // CHECK: 111 0 0 0 + // CHECK: 222 1 0 1 ps++; } @@ -72,7 +72,7 @@ void f1() { printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps[-2].x, s0_hostaddr == &ps[-2].x); // EXPECTED: 333 1 1 0 - // CHECK: 111 1 0 0 + // CHECK: 333 1 0 1 } // The following map(from:ps) should not bring back ps, because ps is an @@ -80,7 +80,8 @@ void f1() { // location, &s[0], on host. #pragma omp target exit data map(always, from : ps) printf("%d %d\n", ps->x, ps == &s[0]); - // CHECK: 111 1 + // EXPECTED: 111 1 + // CHECK: 333 0 #pragma omp target exit data map(delete : ps, s) } diff --git a/offload/test/mapping/map_structptr_and_member_local.c b/offload/test/mapping/map_structptr_and_member_local.c index 9e59551ad3d6c..a9db3eefbc54f 100644 --- a/offload/test/mapping/map_structptr_and_member_local.c +++ b/offload/test/mapping/map_structptr_and_member_local.c @@ -51,7 +51,7 @@ void f1() { printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps->x, s0_hostaddr == &ps->x); // EXPECTED: 111 1 1 0 - // CHECK: 111 0 1 0 + // CHECK: 111 1 0 1 ps++; } @@ -61,7 +61,7 @@ void f1() { printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps[-1].x, s0_hostaddr == &ps[-1].x); // EXPECTED: 222 1 1 0 - // CHECK: 111 0 0 0 + // CHECK: 222 1 0 1 ps++; } @@ -71,7 +71,7 @@ void f1() { printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps[-2].x, s0_hostaddr == &ps[-2].x); // EXPECTED: 333 1 1 0 - // CHECK: 111 1 0 0 + // CHECK: 333 1 0 1 } // The following map(from:ps) should not bring back ps, because ps is an @@ -79,7 +79,8 @@ void f1() { // location, &s[0], on host. #pragma omp target exit data map(always, from : ps) printf("%d %d\n", ps->x, ps == &s[0]); - // CHECK: 111 1 + // EXPECTED: 111 1 + // CHECK: 333 0 #pragma omp target exit data map(delete : ps, s) }