[Clang][OpenMP] Fix segmentation fault when data field is used in is_device_pt.

[jyu2]

Summary

Currently, the field just emit map info for this pointer variable. It is
failed at run time. For the fields, the PartialStruct is created and it
needs call to emitCombinedEntry which create the base that covers all
the pieces.

The change is to generate map info as regular fields.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9074

If VD is nullptr, why do we need this check?

9081

comment with the name of parameter, associated with true value.

[jyu2]

Thanks Alexey! Address the comment.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9074

Yes. You are right. It is unnecessary change. I removed. Thanks. Jennifer

9081

Sorry. Added.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9076

VD is still nullptr here, why do we need to lookup for it?

[jyu2]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9076

VD is nullptr, it represent of this pointer. And it is DevPointersMap. Thanks.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9080

Shall it be OMP_MAP_LITERAL? Previously we emitted OMP_MAP_LITERAL for such items (line 9062)

clang/test/OpenMP/target_is_device_ptr_codegen.cpp

260

// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double**
// CK2-DAG: store double** [[A]], double*** [[CP0]]

?

284–287

Format and checks.

311

same

[jyu2]

Thanks Alexey. This to fix test format.

[ABataev]

clang/test/OpenMP/target_is_device_ptr_codegen.cpp

303

What about this one? Is it used in the test?

[jyu2]

Fix test problem.

clang/test/OpenMP/target_is_device_ptr_codegen.cpp

303

Th B9 is used for size calculate. But I did not check that. So I removed. Thanks. Jennifer

[ABataev]

LG

[alexfh]

It looks like this commit introduces an AddressSanitizer: stack-use-after-scope error:

==2796==ERROR: AddressSanitizer: stack-use-after-scope on address 0x7f544e9acc20 at pc 0x55f1f7f4b83e bp 0x7ffdfb37c560 sp 0x7ffdfb37c558
READ of size 4 at 0x7f544e9acc20 thread T0
    #0 0x55f1f7f4b83d in find<const clang::OpenMPMapModifierKind *, clang::OpenMPMapModifierKind> toolchain/bin/../include/c++/v1/__algorithm/find.h:25:9
    #1 0x55f1f7f4b83d in is_contained<llvm::ArrayRef<clang::OpenMPMapModifierKind> &, clang::OpenMPMapModifierKind> llvm-project/llvm/include/llvm/ADT/STLExtras.h:1684:10
    #2 0x55f1f7f4b83d in (anonymous namespace)::MappableExprsHandler::getMapTypeBits(clang::OpenMPMapClauseKind, llvm::ArrayRef<clang::OpenMPMapModifierKind>, llvm::ArrayRef<clang::OpenMPMotionModifierKind>, bool, bool, bool, bool) const llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:7507:9
    #3 0x55f1f7f460b5 in (anonymous namespace)::MappableExprsHandler::generateInfoForComponentList(clang::OpenMPMapClauseKind, llvm::ArrayRef<clang::OpenMPMapModifierKind>, llvm::ArrayRef<clang::OpenMPMotionModifierKind>, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, (anonymous namespace)::MappableExprsHandler::MapCombinedInfoTy&, (anonymous namespace)::MappableExprsHandler::StructRangeInfoTy&, bool, bool, clang::ValueDecl const*, bool, clang::ValueDecl const*, clang::Expr const*, llvm::ArrayRef<llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>>) const llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8078:45
    #4 0x55f1f7f592f7 in generateInfoForCapture llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:9272:9
    #5 0x55f1f7f592f7 in clang::CodeGen::CGOpenMPRuntime::emitTargetCall(clang::CodeGen::CodeGenFunction&, clang::OMPExecutableDirective const&, llvm::Function*, llvm::Value*, clang::Expr const*, llvm::PointerIntPair<clang::Expr const*, 2u, clang::OpenMPDeviceClauseModifier, llvm::PointerLikeTypeTraits<clang::Expr const*>, llvm::PointerIntPairInfo<clang::Expr const*, 2u, llvm::PointerLikeTypeTraits<clang::Expr const*>>>, llvm::function_ref<llvm::Value* (clang::CodeGen::CodeGenFunction&, clang::OMPLoopDirective const&)>)::$_22::operator()(clang::CodeGen::CodeGenFunction&, clang::CodeGen::PrePostActionTy&) const llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:10411:19
    #6 0x55f1f7ec9f8b in clang::CodeGen::RegionCodeGenTy::operator()(clang::CodeGen::CodeGenFunction&) const llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:603:5
    #7 0x55f1f7f20fb5 in clang::CodeGen::CGOpenMPRuntime::emitTargetCall(clang::CodeGen::CodeGenFunction&, clang::OMPExecutableDirective const&, llvm::Function*, llvm::Value*, clang::Expr const*, llvm::PointerIntPair<clang::Expr const*, 2u, clang::OpenMPDeviceClauseModifier, llvm::PointerLikeTypeTraits<clang::Expr const*>, llvm::PointerIntPairInfo<clang::Expr const*, 2u, llvm::PointerLikeTypeTraits<clang::Expr const*>>>, llvm::function_ref<llvm::Value* (clang::CodeGen::CodeGenFunction&, clang::OMPLoopDirective const&)>) llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:10497:7
    #8 0x55f1f7dbbad3 in emitCommonOMPTargetDirective(clang::CodeGen::CodeGenFunction&, clang::OMPExecutableDirective const&, clang::CodeGen::RegionCodeGenTy const&) llvm-project/clang/lib/CodeGen/CGStmtOpenMP.cpp:6639:26
    #9 0x55f1f7ddb476 in clang::CodeGen::CodeGenFunction::EmitOMPTargetDirective(clang::OMPTargetDirective const&) llvm-project/clang/lib/CodeGen/CGStmtOpenMP.cpp:6675:3
    #10 0x55f1f7d877db in clang::CodeGen::CodeGenFunction::EmitCompoundStmtWithoutScope(clang::CompoundStmt const&, bool, clang::CodeGen::AggValueSlot) llvm-project/clang/lib/CodeGen/CGStmt.cpp:531:7
    #11 0x55f1f822ede6 in clang::CodeGen::CodeGenFunction::EmitFunctionBody(clang::Stmt const*) llvm-project/clang/lib/CodeGen/CodeGenFunction.cpp:1234:5
    #12 0x55f1f82304de in clang::CodeGen::CodeGenFunction::GenerateCode(clang::GlobalDecl, llvm::Function*, clang::CodeGen::CGFunctionInfo const&) llvm-project/clang/lib/CodeGen/CodeGenFunction.cpp:1442:5
    #13 0x55f1f826ebdf in clang::CodeGen::CodeGenModule::EmitGlobalFunctionDefinition(clang::GlobalDecl, llvm::GlobalValue*) llvm-project/clang/lib/CodeGen/CodeGenModule.cpp:5243:26
    #14 0x55f1f8261f55 in clang::CodeGen::CodeGenModule::EmitGlobalDefinition(clang::GlobalDecl, llvm::GlobalValue*) llvm-project/clang/lib/CodeGen/CodeGenModule.cpp:3510:9
    #15 0x55f1f824b75d in clang::CodeGen::CodeGenModule::EmitDeferred() llvm-project/clang/lib/CodeGen/CodeGenModule.cpp:2681:5
    #16 0x55f1f824b79d in clang::CodeGen::CodeGenModule::EmitDeferred() llvm-project/clang/lib/CodeGen/CodeGenModule.cpp:2687:7
    #17 0x55f1f824742e in clang::CodeGen::CodeGenModule::Release() llvm-project/clang/lib/CodeGen/CodeGenModule.cpp:526:3
    #18 0x55f1f848b03f in (anonymous namespace)::CodeGeneratorImpl::HandleTranslationUnit(clang::ASTContext&) llvm-project/clang/lib/CodeGen/ModuleBuilder.cpp:286:18
    #19 0x55f1f7b5cf7a in clang::BackendConsumer::HandleTranslationUnit(clang::ASTContext&) llvm-project/clang/lib/CodeGen/CodeGenAction.cpp:306:14
    #20 0x55f1f96d724a in clang::ParseAST(clang::Sema&, bool, bool) llvm-project/clang/lib/Parse/ParseAST.cpp:198:13
    #21 0x55f1f7b5816f in clang::CodeGenAction::ExecuteAction() llvm-project/clang/lib/CodeGen/CodeGenAction.cpp:1144:30
    #22 0x55f1f923f912 in clang::FrontendAction::Execute() llvm-project/clang/lib/Frontend/FrontendAction.cpp:1037:8
    #23 0x55f1f916e16d in clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) llvm-project/clang/lib/Frontend/CompilerInstance.cpp:1035:33
    #24 0x55f1f742275c in clang::ExecuteCompilerInvocation(clang::CompilerInstance*) llvm-project/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp:266:25
    #25 0x55f1f741e5bd in cc1_main(llvm::ArrayRef<char const*>, char const*, void*) llvm-project/clang/tools/driver/cc1_main.cpp:250:15
    #26 0x55f1f7409014 in ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) llvm-project/clang/tools/driver/driver.cpp:317:12
    #27 0x55f1f74087e2 in clang_main(int, char**) llvm-project/clang/tools/driver/driver.cpp:388:12

Address 0x7f544e9acc20 is located in stack of thread T0 at offset 3104 in frame
    #0 0x55f1f7f5512f in clang::CodeGen::CGOpenMPRuntime::emitTargetCall(clang::CodeGen::CodeGenFunction&, clang::OMPExecutableDirective const&, llvm::Function*, llvm::Value*, clang::Expr const*, llvm::PointerIntPair<clang::Expr const*, 2u, clang::OpenMPDeviceClauseModifier, llvm::PointerLikeTypeTraits<clang::Expr const*>, llvm::PointerIntPairInfo<clang::Expr const*, 2u, llvm::PointerLikeTypeTraits<clang::Expr const*>>>, llvm::function_ref<llvm::Value* (clang::CodeGen::CodeGenFunction&, clang::OMPLoopDirective const&)>)::$_22::operator()(clang::CodeGen::CodeGenFunction&, clang::CodeGen::PrePostActionTy&) const llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:10378

  This frame has 116 object(s):
    [32, 40) 'ConstFoundBucket.i.i800.i'
    [64, 72) 'retval.i.i732.i'
    [96, 104) 'ConstFoundBucket.i.i635.i'
    [128, 136) 'ConstFoundBucket.i.i606.i'
    [160, 168) 'Val.addr.i.i337'
    [192, 200) 'ConstFoundBucket.i.i.i338'
    [224, 264) 'agg.tmp7.i.i'
    [304, 312) 'retval.i.i.i339'
    [336, 352) 'VDAddr.i' (line 8948)
    [368, 472) 'VDLVal.i' (line 8950)
    [512, 536) 'Captures.i' (line 8951)
    [576, 584) 'ThisCapture.i' (line 8952)
    [608, 712) 'ThisLVal.i' (line 8955)
    [752, 856) 'ThisLValVal.i' (line 8957)
    [896, 904) 'ref.tmp55.i' (line 8958)
    [928, 968) 'ref.tmp111.i' (line 8963)
    [1008, 1112) 'VarLVal.i' (line 8978)
    [1152, 1256) 'VarLValVal.i' (line 8980)
    [1296, 1304) 'ref.tmp186.i340' (line 8981)
    [1328, 1368) 'ref.tmp253.i' (line 8986)
    [1408, 1432) 'VarRVal.i' (line 8991)
    [1472, 1480) 'ref.tmp267.i' (line 8992)
    [1504, 1512) 'TheBucket.i667.i'
    [1536, 1544) 'Val.addr.i.i.i'
    [1568, 1576) 'TheBucket.i85.i.i'
    [1600, 1608) 'TheBucket.i.i.i'
    [1632, 1640) 'ref.tmp.i.i' (line 8348)
    [1664, 1672) 'TheBucket.i.i242'
    [1696, 1736) 'ref.tmp36.i' (line 9292)
    [1776, 1816) 'ref.tmp101.i' (line 9306)
    [1856, 1864) 'ref.tmp123.i' (line 9314)
    [1888, 1928) 'ref.tmp179.i' (line 9321)
    [1968, 1976) 'ref.tmp193.i' (line 9328)
    [2000, 2104) 'agg.tmp224.i' (line 9351)
    [2144, 2152) 'ConstFoundBucket.i.i.i216'
    [2176, 2184) 'ConstFoundBucket.i.i.i205'
    [2208, 2216) 'ConstFoundBucket.i.i1266.i'
    [2240, 2248) 'Val.addr.i1267.i'
    [2272, 2280) '__comp.i.i1185.i'
    [2304, 2312) 'retval.i.i.i'
    [2336, 2344) 'ConstFoundBucket.i.i.i'
    [2368, 2369) '__comp.i.i.i'
    [2384, 2392) 'Val.addr.i850.i'
    [2416, 2424) 'TheBucket.i851.i'
    [2448, 2456) 'Val.addr.i778.i'
    [2480, 2488) 'TheBucket.i779.i'
    [2512, 2520) 'Val.addr.i.i162'
    [2544, 2552) 'TheBucket.i.i163'
    [2576, 2584) 'Arg.addr.i' (line 9278)
    [2608, 2616) 'VD.i' (line 9043)
    [2640, 2680) 'ref.tmp.i' (line 9059)
    [2720, 2992) 'DeclComponentLists.i' (line 9073)
    [3056, 3072) 'MCL.i' (line 9078)
    [3088, 3092) 'ref.tmp94.i' (line 9079)
    [3104, 3108) 'ref.tmp95.i' (line 9079) <== Memory access at offset 3104 is inside this variable
    [3120, 3121) 'ref.tmp96.i' (line 9079)
    [3136, 3144) 'ref.tmp97.i' (line 9079)
    [3168, 3176) 'ref.tmp98.i' (line 9079)
    [3200, 3216) '__begin2.i' (line 9085)
    [3232, 3376) 'ref.tmp135.i' (line 9087)
    [3440, 3512) '__begin3138.i' (line 9087)
    [3552, 3584) 'L.i' (line 9087)
    [3616, 3624) 'Mapper.i' (line 9088)
    [3648, 3656) 'E.i' (line 9090)
    [3680, 3696) 'Components.i' (line 9091)
    [3712, 3716) 'ref.tmp178.i' (line 9096)
    [3728, 3744) 'ref.tmp182.i' (line 9096)
    [3760, 3761) 'ref.tmp186.i' (line 9096)
    [3776, 4136) 'OverlappedData.i' (line 9118)
    [4208, 4224) 'Components215.i' (line 9126)
    [4240, 4244) 'MapType.i' (line 9127)
    [4256, 4272) 'MapModifiers.i' (line 9128)
    [4288, 4289) 'IsImplicit.i' (line 9129)
    [4304, 4312) 'Mapper216.i' (line 9130)
    [4336, 4344) 'VarRef.i' (line 9131)
    [4368, 4416) 'ref.tmp219.i' (line 9132)
    [4448, 4464) 'Components1.i' (line 9136)
    [4480, 4528) 'ref.tmp242.i' (line 9137)
    [4560, 4568) 'ref.tmp409.i' (line 9176)
    [4592, 4640) 'Layout.i' (line 9182)
    [4672, 4688) 'Components544.i' (line 9244)
    [4704, 4708) 'MapType545.i' (line 9245)
    [4720, 4736) 'MapModifiers546.i' (line 9246)
    [4752, 4753) 'IsImplicit547.i' (line 9247)
    [4768, 4776) 'Mapper548.i' (line 9248)
    [4800, 4808) 'VarRef549.i' (line 9249)
    [4832, 4880) 'ref.tmp552.i' (line 9250)
    [4912, 4928) 'agg.tmp569.i' (line 9278)
    [4944, 4960) 'Components592.i' (line 9262)
    [4976, 4980) 'MapType593.i' (line 9263)
    [4992, 5008) 'MapModifiers594.i' (line 9264)
    [5024, 5025) 'IsImplicit595.i' (line 9265)
    [5040, 5048) 'Mapper596.i' (line 9266)
    [5072, 5080) 'VarRef597.i' (line 9267)
    [5104, 5152) 'ref.tmp600.i' (line 9268)
    [5184, 5200) 'agg.tmp621.i' (line 9278)
    [5216, 5224) 'Val.addr.i.i'
    [5248, 5256) 'TheBucket.i.i'
    [5280, 6312) 'CombinedInfo' (line 10380)
    [6448, 6536) 'MEHandler' (line 10383)
    [6576, 6600) 'LambdaPointers' (line 10384)
    [6640, 6664) 'MappedVarSet' (line 10385)
    [6704, 6712) 'RI' (line 10387)
    [6736, 7768) 'CurInfo' (line 10392)
    [7904, 9024) 'PartialStruct' (line 10393)
    [9152, 9192) 'ref.tmp' (line 10401)
    [9232, 9240) 'ref.tmp56' (line 10413)
    [9264, 9272) 'ref.tmp63' (line 10415)
    [9296, 9392) 'Info' (line 10453)
    [9424, 9425) 'ref.tmp192' (line 10456)
    [9440, 9456) 'ref.tmp193' (line 10462)
    [9472, 9488) 'ref.tmp210' (line 10464)
    [9504, 9520) 'ref.tmp229' (line 10466)
    [9536, 9552) 'ref.tmp246' (line 10468)
    [9568, 9592) 'ref.tmp271' (line 10473)
    [9632, 9656) 'ref.tmp280' (line 10475)


SUMMARY: AddressSanitizer: stack-use-after-scope toolchain/bin/../include/c++/v1/__algorithm/find.h:25:9 in find<const clang::OpenMPMapModifierKind *, clang::OpenMPMapModifierKind>
Shadow bytes around the buggy address:
  0x0feb09d2d930: f2 f2 f8 f2 f2 f2 f8 f2 f2 f2 f8 f2 f2 f2 f8 f2
  0x0feb09d2d940: f2 f2 00 f2 f2 f2 00 f2 f2 f2 f8 f8 f8 f8 f8 f2
  0x0feb09d2d950: f2 f2 f2 f2 00 00 00 00 00 00 00 00 00 00 00 00
  0x0feb09d2d960: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
  0x0feb09d2d970: 00 00 00 00 00 00 f2 f2 f2 f2 f2 f2 f2 f2 f8 f8
=>0x0feb09d2d980: f2 f2 f8 f2[f8]f2 f8 f2 f8 f2 f2 f2 f8 f2 f2 f2
  0x0feb09d2d990: f8 f8 f2 f2 f8 f8 f8 f8 f8 f8 f8 f8 f8 f8 f8 f8
  0x0feb09d2d9a0: f8 f8 f8 f8 f8 f8 f2 f2 f2 f2 f2 f2 f2 f2 f8 f8
  0x0feb09d2d9b0: f8 f8 f8 f8 f8 f8 f8 f2 f2 f2 f2 f2 f8 f8 f8 f8
  0x0feb09d2d9c0: f2 f2 f2 f2 f8 f2 f2 f2 f8 f2 f2 f2 f8 f8 f2 f2
  0x0feb09d2d9d0: f8 f2 f8 f8 f2 f2 f8 f2 00 00 00 00 00 00 00 00
Shadow byte legend (one shadow byte represents 8 application bytes):
  Addressable:           00
  Partially addressable: 01 02 03 04 05 06 07 
  Heap left redzone:       fa
  Freed heap region:       fd
  Stack left redzone:      f1
  Stack mid redzone:       f2
  Stack right redzone:     f3
  Stack after return:      f5
  Stack use after scope:   f8
  Global redzone:          f9
  Global init order:       f6
  Poisoned by user:        f7
  Container overflow:      fc
  Array cookie:            ac
  Intra object redzone:    bb
  ASan internal:           fe
  Left alloca redzone:     ca
  Right alloca redzone:    cb
==2796==ABORTING
FileCheck error: '<stdin>' is empty.
FileCheck command line:  llvm-project/llvm/FileCheck --allow-unused-prefixes llvm-project/clang/test/OpenMP/target_is_device_ptr_codegen.cpp --check-prefix CK2 --check-prefix CK2-64

[jyu2]

Hi @alexfh,

How could I reproduce the problem? Thanks.
Thanks.

Jennifer

[alexfh]

In D129608#3735496, @jyu2 wrote:

Hi @alexfh,

How could I reproduce the problem? Thanks.
Thanks.

Jennifer

It should be reproducible by running the tests with the address sanitizer enabled. But it looks like this might have already been fixed in f385eaf48f5c4123b61272cd243c339d6f8526ed.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

void generateInfoForCapture(const CapturedStmt::Capture *Cap,
     // for map(to: lambda): skip here, processing it in
     // generateDefaultMapInfo
     if (LambdasMap.count(VD))
       return;
 
     // If this declaration appears in a is_device_ptr clause we just have to
     // pass the pointer by value. If it is a reference to a declaration, we just
     // pass its value.
-    if (DevPointersMap.count(VD)) {
+    if (VD && DevPointersMap.count(VD)) {
       CombinedInfo.Exprs.push_back(VD);
       CombinedInfo.BasePointers.emplace_back(Arg, VD);
       CombinedInfo.Pointers.push_back(Arg);
       CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
           CGF.getTypeSize(CGF.getContext().VoidPtrTy), CGF.Int64Ty,
           /*isSigned=*/true));
       CombinedInfo.Types.push_back(
           (Cap->capturesVariable() ? OMP_MAP_TO : OMP_MAP_LITERAL) |
           OMP_MAP_TARGET_PARAM);
       CombinedInfo.Mappers.push_back(nullptr);
       return;
     }
 
     using MapData =
         std::tuple<OMPClauseMappableExprCommon::MappableExprComponentListRef,
                    OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool,
                    const ValueDecl *, const Expr *>;
     SmallVector<MapData, 4> DeclComponentLists;
+    // For member fields list in is_device_ptr, store it in
+    // DeclComponentLists for generating components info.
+    auto It = DevPointersMap.find(VD);
+    if (It != DevPointersMap.end())
+      for (const auto MCL : It->second)
+        DeclComponentLists.emplace_back(
+            MCL, OMPC_MAP_to, OMPC_MAP_MODIFIER_unknown, /*IsImpicit = */ true,
+            nullptr, nullptr);
     assert(CurDir.is<const OMPExecutableDirective *>() &&
            "Expect a executable directive");
     const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
     for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
       const auto *EI = C->getVarRefs().begin();
       for (const auto L : C->decl_component_lists(VD)) {
         const ValueDecl *VDecl, *Mapper;
         // The Expression is not correct if the mapping is implicit

clang/test/OpenMP/target_is_device_ptr_codegen.cpp

 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
+// CK2-DAG: [[A:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1:%.+]], i32 0, i32 0
 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
-// CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
-// CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
+// CK2-DAG: store [[ST]]* [[THIS1]], [[ST]]** [[CBP0]]
+// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double**
+// CK2-DAG: store double** [[A]], double*** [[CP0]]
 #pragma omp target is_device_ptr(a)
     {
       a++;
     }
 
 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
+// CK2-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CK2-DAG: store i64* [[SIZE:%.+]], i64** [[SARG]]
 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
+// CK2-DAG: [[S:%[^,]+]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK2-DAG: [[SIZE:%[^,]+]] = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i32 0, i32 0
+// CK2-DAG: store i64 [[S]], i64* [[SIZE]]
+// CK2-DAG: [[B:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
-// CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
-// CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
+// CK2-DAG:  store %struct.ST* [[THIS1]], %struct.ST** [[CBP0]]
+// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
+// CK2-DAG: store double*** [[B]], double**** [[CP0]]
 #pragma omp target is_device_ptr(b)
     {
       b++;
     }
 
 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
+// CK2-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CK2-DAG: store i64* [[SIZE:%.+]], i64** [[SARG]]
 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
+// CK2-DAG: [[A8:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 0
+// CK2-DAG: [[B9:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
+// CK2-DAG: [[S:%[^,]+]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK2-DAG: store i64 [[S]], i64* [[SIZE:%.+]]
+
 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
-// CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
-// CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
+// CK2-DAG:  store %struct.ST* [[THIS1]], %struct.ST** [[CBP0]]
+// CH2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to to double***
+// CK2-DAG: store double** [[A8]], double*** [[TMP64:%.+]]
 #pragma omp target is_device_ptr(a, b)
     {
       a++;
       b++;
     }
   }
 };
 

openmp/libomptarget/test/mapping/is_device_ptr.cpp

+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <assert.h>
+#include <iostream>
+#include <omp.h>
+
+struct view {
+  const int size = 10;
+  int *data_host;
+  int *data_device;
+  void foo() {
+    std::size_t bytes = size * sizeof(int);
+    const int host_id = omp_get_initial_device();
+    const int device_id = omp_get_default_device();
+    data_host = (int *)malloc(bytes);
+    data_device = (int *)omp_target_alloc(bytes, device_id);
+#pragma omp target teams distribute parallel for is_device_ptr(data_device)
+    for (int i = 0; i < size; ++i)
+      data_device[i] = i;
+    omp_target_memcpy(data_host, data_device, bytes, 0, 0, host_id, device_id);
+    for (int i = 0; i < size; ++i)
+      assert(data_host[i] == i);
+  }
+};
+
+int main() {
+  view a;
+  a.foo();
+  // CHECK: PASSED
+  printf("PASSED\n");
+}