llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-clang Author: Abhinav Gaba (abhinavgaba) <details> <summary>Changes</summary> For the following: ```c int *p; #pragma omp target map(p[0]) // (A) (void)p; #pragma omp target map(p) // (B) (void)p; #pragma omp target map(p, p[0]) // (C) (void)p; #pragma omp target map(p[0], p) // (D) (void)p; ``` For (A), the pointer `p` is predetermined `firstprivate`, so it should be (and is) captured by-copy. However, for (B), (C), and (D), since `p` is already listed in a `map` clause, it's not predetermined `firstprivate`, and hence, should be captured by-reference, like any other mapped variable. To ensure the correct handling of (C) and (D), the following changes were made: 1. In SemaOpenMP, we now ensure that `p` is marked to be captured by-reference in these cases. 2. We no longer ignore `map(p)` during codegen of `target` constructs, even if there's another map like `map(p[0])` that would have been mapped using a PTR_AND_OBJ map. 3. For cases like (D), we now handle `map(p)` before `map(p[0])`, so the former gets the TARGET_PARAM flag and sets the kernel argument. --- Patch is 34.47 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/145454.diff 5 Files Affected: - (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+33-4) - (modified) clang/lib/Sema/SemaOpenMP.cpp (+42-9) - (modified) clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp (+147-143) - (added) offload/test/mapping/map_ptr_and_subscript_global.c (+66) - (added) offload/test/mapping/map_ptr_and_subscript_local.c (+66) ``````````diff diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 8ccc37ef98a74..39571105e26b2 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7270,8 +7270,14 @@ class MappableExprsHandler { // of arguments, hence MEMBER_OF(4) // // map(p, p[:100]) + // For "pragma omp target": + // &p, &p, sizeof(p), TARGET_PARAM | TO | FROM + // &p, &p[0], 100*sizeof(float), PTR_AND_OBJ | TO | FROM (*) + // Otherwise: // ===> map(p[:100]) // &p, &p[0], 100*sizeof(float), TARGET_PARAM | PTR_AND_OBJ | TO | FROM + // (*) We need to use PTR_AND_OBJ here to ensure that the mapped copies of + // p and p[0] get attached. // Track if the map information being generated is the first for a capture. bool IsCaptureFirstInfo = IsFirstComponentList; @@ -7289,14 +7295,26 @@ class MappableExprsHandler { // components. bool IsExpressionFirstInfo = true; bool FirstPointerInComplexData = false; + bool SkipStandalonePtrMapping = false; Address BP = Address::invalid(); const Expr *AssocExpr = I->getAssociatedExpression(); const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr); const auto *OASE = dyn_cast<ArraySectionExpr>(AssocExpr); const auto *OAShE = dyn_cast<OMPArrayShapingExpr>(AssocExpr); - if (AreBothBasePtrAndPteeMapped && std::next(I) == CE) + // For map(p, p[0]) on a "target" construct, we need to map "p" by itself + // as it has to be passed by-reference as the kernel argument. + // For other constructs, we can skip mapping "p" because the PTR_AND_OBJ + // mapping for map(p[0]) will take care of mapping p as well. + SkipStandalonePtrMapping = + AreBothBasePtrAndPteeMapped && + (!isa<const OMPExecutableDirective *>(CurDir) || + !isOpenMPTargetExecutionDirective( + cast<const OMPExecutableDirective *>(CurDir)->getDirectiveKind())); + + if (SkipStandalonePtrMapping && std::next(I) == CE) return; + if (isa<MemberExpr>(AssocExpr)) { // The base is the 'this' pointer. The content of the pointer is going // to be the base of the field being mapped. @@ -7672,7 +7690,7 @@ class MappableExprsHandler { getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit, !IsExpressionFirstInfo || RequiresReference || FirstPointerInComplexData || IsMemberReference, - AreBothBasePtrAndPteeMapped || + SkipStandalonePtrMapping || (IsCaptureFirstInfo && !RequiresReference), IsNonContiguous); @@ -8811,8 +8829,19 @@ class MappableExprsHandler { ++EI; } } - llvm::stable_sort(DeclComponentLists, [](const MapData &LHS, - const MapData &RHS) { + llvm::stable_sort(DeclComponentLists, [VD](const MapData &LHS, + const MapData &RHS) { + // For cases like map(p, p[0], p[0][0]), the shortest map, like map(p) + // in this case, should be handled first, to ensure that it gets the + // TARGET_PARAM flag. + OMPClauseMappableExprCommon::MappableExprComponentListRef Components = + std::get<0>(LHS); + OMPClauseMappableExprCommon::MappableExprComponentListRef ComponentsR = + std::get<0>(RHS); + if (VD && VD->getType()->isAnyPointerType() && Components.size() == 1 && + ComponentsR.size() > 1) + return true; + ArrayRef<OpenMPMapModifierKind> MapModifiers = std::get<2>(LHS); OpenMPMapClauseKind MapType = std::get<1>(RHS); bool HasPresent = diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 00f4658180807..02e4e7b910d2e 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -2146,6 +2146,7 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level, // | ptr | n.a. | - | x | - | - | bycopy| // | ptr | n.a. | x | - | - | - | null | // | ptr | n.a. | - | - | - | x | byref | + // | ptr | n.a. | - | - | - | x, x[] | bycopy| // | ptr | n.a. | - | - | - | x[] | bycopy| // | ptr | n.a. | - | - | x | | bycopy| // | ptr | n.a. | - | - | x | x | bycopy| @@ -2171,18 +2172,22 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level, // - For pointers mapped by value that have either an implicit map or an // array section, the runtime library may pass the NULL value to the // device instead of the value passed to it by the compiler. + // - If both a pointer an a dereference of it are mapped, then the pointer + // should be passed by reference. if (Ty->isReferenceType()) Ty = Ty->castAs<ReferenceType>()->getPointeeType(); - // Locate map clauses and see if the variable being captured is referred to - // in any of those clauses. Here we only care about variables, not fields, - // because fields are part of aggregates. + // Locate map clauses and see if the variable being captured is mapped by + // itself, or referred to, in any of those clauses. Here we only care about + // variables, not fields, because fields are part of aggregates. bool IsVariableAssociatedWithSection = false; + bool IsVariableItselfMapped = false; DSAStack->checkMappableExprComponentListsForDeclAtLevel( D, Level, [&IsVariableUsedInMapClause, &IsVariableAssociatedWithSection, + &IsVariableItselfMapped, D](OMPClauseMappableExprCommon::MappableExprComponentListRef MapExprComponents, OpenMPClauseKind WhereFoundClauseKind) { @@ -2198,8 +2203,19 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level, assert(EI != EE && "Invalid map expression!"); - if (isa<DeclRefExpr>(EI->getAssociatedExpression())) - IsVariableUsedInMapClause |= EI->getAssociatedDeclaration() == D; + if (isa<DeclRefExpr>(EI->getAssociatedExpression()) && + EI->getAssociatedDeclaration() == D) { + IsVariableUsedInMapClause = true; + + // If the component list has only one element, it's for mapping the + // variable itself, like map(p). This takes precedence in + // determining how it's captured, so we don't need to look further + // for any other maps that use the variable (like map(p[0]) etc.) + if (MapExprComponents.size() == 1) { + IsVariableItselfMapped = true; + return true; + } + } ++EI; if (EI == EE) @@ -2213,8 +2229,10 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level, isa<MemberExpr>(EI->getAssociatedExpression()) || isa<OMPArrayShapingExpr>(Last->getAssociatedExpression())) { IsVariableAssociatedWithSection = true; - // There is nothing more we need to know about this variable. - return true; + // We've found a case like map(p[0]) or map(p->a) or map(*p), + // so we are done with this particular map, but we need to keep + // looking in case we find a map(p). + return false; } // Keep looking for more map info. @@ -2223,8 +2241,23 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level, if (IsVariableUsedInMapClause) { // If variable is identified in a map clause it is always captured by - // reference except if it is a pointer that is dereferenced somehow. - IsByRef = !(Ty->isPointerType() && IsVariableAssociatedWithSection); + // reference except if it is a pointer that is dereferenced somehow, but + // not itself mapped. + // + // OpenMP 6.0, 7.1.1: Data sharing attribute rules, variables referenced + // in a construct:: + // If a list item in a has_device_addr clause or in a map clause on the + // target construct has a base pointer, and the base pointer is a scalar + // variable *that is not a list item in a map clause on the construct*, + // the base pointer is firstprivate. + // + // OpenMP 4.5, 2.15.1.1: Data-sharing Attribute Rules for Variables + // Referenced in a Construct: + // If an array section is a list item in a map clause on the target + // construct and the array section is derived from a variable for which + // the type is pointer then that variable is firstprivate. + IsByRef = IsVariableItselfMapped || + !(Ty->isPointerType() && IsVariableAssociatedWithSection); } else { // By default, all the data that has a scalar type is mapped by copy // (except for reduction variables). diff --git a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp index 87fa7fe462daa..9a8f234da718c 100644 --- a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp +++ b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp @@ -1,174 +1,178 @@ -// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ // RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s // expected-no-diagnostics #ifndef HEADER #define HEADER +// CHECK: @.[[KERNEL00:__omp_offloading_.*foov_l[0-9]+]].region_id = weak constant i8 0 +// CHECK: [[SIZE00:@.+]] = private unnamed_addr constant [2 x i64] [i64 {{8|4}}, i64 8] +// CHECK: [[MYTYPE00:@.+]] = private unnamed_addr constant [2 x i64] [i64 35, i64 19] + +// CHECK: @.[[KERNEL01:__omp_offloading_.*foov_l[0-9]+]].region_id = weak constant i8 0 +// CHECK: [[SIZE01:@.+]] = private unnamed_addr constant [2 x i64] [i64 {{8|4}}, i64 4] +// CHECK: [[MYTYPE01:@.+]] = private unnamed_addr constant [2 x i64] [i64 35, i64 19] + +// CHECK: @.[[KERNEL02:__omp_offloading_.*foov_l[0-9]+]].region_id = weak constant i8 0 +// CHECK: [[SIZE02:@.+]] = private unnamed_addr constant [2 x i64] [i64 {{8|4}}, i64 4] +// CHECK: [[MYTYPE02:@.+]] = private unnamed_addr constant [2 x i64] [i64 35, i64 19] + +// CHECK: [[SIZE03:@.+]] = private unnamed_addr constant [1 x i64] [i64 4] +// CHECK: [[MYTYPE03:@.+]] = private unnamed_addr constant [1 x i64] [i64 51] + extern void *malloc (int __size) throw () __attribute__ ((__malloc__)); +// CHECK-LABEL: define{{.*}}@_Z3foov{{.*}}( void foo() { int *ptr = (int *) malloc(3 * sizeof(int)); +// Region 00 +// &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM +// &ptr, &ptr[0], 2 * sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ +// +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.[[KERNEL00]].region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 +// CHECK-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] +// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 +// CHECK-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] +// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] +// +// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]] +// CHECK-DAG: store ptr [[VAR0]], ptr [[P0]] +// +// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP1]] +// CHECK-DAG: store ptr [[RVAR00:%.+]], ptr [[P1]] +// +// CHECK-DAG: [[RVAR00]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}} 0 +// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]] +// +// CHECK-DAG: call void @[[KERNEL00]](ptr [[VAR0]]) #pragma omp target map(ptr, ptr[0:2]) { ptr[1] = 6; } + +// Region 01 +// &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM +// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ +// +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.[[KERNEL01]].region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 +// CHECK-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] +// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 +// CHECK-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] +// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] +// +// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]] +// CHECK-DAG: store ptr [[VAR0]], ptr [[P0]] +// +// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP1]] +// CHECK-DAG: store ptr [[RVAR02:%.+]], ptr [[P1]] +// +// CHECK-DAG: [[RVAR02]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}} 2 +// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]] +// +// CHECK-DAG: call void @[[KERNEL01]](ptr [[VAR0]]) #pragma omp target map(ptr, ptr[2]) { ptr[2] = 8; } - #pragma omp target data map(ptr, ptr[2]) + +// Region 02 +// &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM +// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PTR_AND_OBJ +// +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.[[KERNEL02]].region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 +// CHECK-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] +// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 +// CHECK-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] +// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] +// +// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]] +// CHECK-DAG: store ptr [[VAR0]], ptr [[P0]] +// +// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP1]] +// CHECK-DAG: store ptr [[RVAR02:%.+]], ptr [[P1]] +// +// CHECK-DAG: [[RVAR02]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}} 2 +// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]] +// +// CHECK-DAG: call void @[[KERNEL02]](ptr [[VAR0]]) + #pragma omp target map(ptr[2], ptr) { ptr[2] = 9; } + +// Region 03 +// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ +// FIXME: PARAM seems to be redundant here. +// +// CHECK-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BPGEP:.+]], ptr [[PGEP:.+]], ptr [[SIZE03]], ptr [[MYTYPE03]], ptr null, ptr null) +// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] +// +// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CHECK-DAG: store ptr [[VAR0:%ptr]], ptr [[BP0]] +// CHECK-DAG: store ptr [[RVAR02:%.+]], ptr [[P0]] +// +// CHECK-DAG: [[RVAR02]] = getelementptr inbounds {{.*}}[[RVAR0:%.+]], i{{.+}} 2 +// CHECK-DAG: [[RVAR0]] = load ptr, ptr [[VAR0]] + #pragma omp target data map(ptr, ptr[2]) + { + ptr[2] = 10; + } } -#endif -// CHECK-LABEL: define {{[^@]+}}@_Z3foov -// CHECK-SAME: () #[[ATTR0:[0-9]+]] { + +// CHECK-LABEL: define internal void +// CHECK-SAME: @[[KERNEL00]](ptr {{[^,]*}}[[PTR:%[^,]+]]) // CHECK-NEXT: entry: -// CHECK-NEXT: [[PTR:%.*]] = alloca ptr, align 8 -// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [1 x ptr], align 8 -// CHECK-NEXT: [[CALL:%.*]] = call noalias noundef ptr @_Z6malloci(i32 noundef signext 12) #[[ATTR3:[0-9]+]] -// CHECK-NEXT: store ptr [[CALL]], ptr [[PTR]], align 8 -// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8 -// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 0 -// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP2]], align 8 -// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 -// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 -// CHECK-NEXT: store ptr null, ptr [[TMP4]], align 8 -// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 -// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 -// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 -// CHECK-NEXT: store i32 3, ptr [[TMP7]], align 4 -// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 -// CHECK-NEXT: store i32 1, ptr [[TMP8]], align 4 -// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 -// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8 -// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 -// CHECK-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8 -// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 -// CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8 -// CHECK-NEXT: [[TMP12:%.*]] = getelement... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/145454 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
