This is an archive of the discontinued LLVM Phabricator instance.

Differential D144873

[OpenMP] Ignore implicit casts on assertion for `use_device_ptr`
ClosedPublic

Authored by jhuber6 on Feb 27 2023, 6:54 AM.

Details

Reviewers
jdoerfert
tianshilei1992
ABataev
ronlieb
carlo.bertolli
Commits
rG853d4059135b: [OpenMP] Ignore implicit casts on assertion for `use_device_ptr`
Summary

There was an assertion triggering when invoking a captured member whose
initializer was in a blase class. This patch fixes it by allowing the
assertion on implicit casts to the base class rather than only the base
class itself.

Fixes https://github.com/llvm/llvm-project/issues/61027

Diff Detail

Event Timeline

jhuber6 created this revision.Feb 27 2023, 6:54 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 27 2023, 6:54 AM
Herald added subscribers: guansong, yaxunl. · View Herald Transcript
jhuber6 requested review of this revision.Feb 27 2023, 6:54 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 27 2023, 6:54 AM
Herald added subscribers: cfe-commits, sstefan1. · View Herald Transcript
tianshilei1992 accepted this revision.Feb 27 2023, 7:32 AM

LG

This revision is now accepted and ready to land.Feb 27 2023, 7:32 AM
Harbormaster completed remote builds in B216213: Diff 500775.Feb 27 2023, 8:13 AM
This revision was landed with ongoing or failed builds.Feb 27 2023, 8:48 AM
Closed by commit rG853d4059135b: [OpenMP] Ignore implicit casts on assertion for `use_device_ptr` (authored by jhuber6). · Explain Why
This revision was automatically updated to reflect the committed changes.
jhuber6 added a commit: rG853d4059135b: [OpenMP] Ignore implicit casts on assertion for `use_device_ptr`.

Revision Contents

PathSize
clang/
lib/
CodeGen/
2 lines
test/
OpenMP/
266 lines

Diff 500815

clang/lib/CodeGen/CGStmtOpenMP.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 7,140 Lines • ▼ Show 20 Lines for (const Expr *PvtVarIt : C.private_copies()) {
// In order to identify the right initializer we need to match the // In order to identify the right initializer we need to match the
// declaration used by the mapping logic. In some cases we may get // declaration used by the mapping logic. In some cases we may get
// OMPCapturedExprDecl that refers to the original declaration. // OMPCapturedExprDecl that refers to the original declaration.
const ValueDecl *MatchingVD = OrigVD; const ValueDecl *MatchingVD = OrigVD;
if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(MatchingVD)) { if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(MatchingVD)) {
// OMPCapturedExprDecl are used to privative fields of the current // OMPCapturedExprDecl are used to privative fields of the current
// structure. // structure.
const auto *ME = cast<MemberExpr>(OED->getInit()); const auto *ME = cast<MemberExpr>(OED->getInit());
assert(isa<CXXThisExpr>(ME->getBase()) && assert(isa<CXXThisExpr>(ME->getBase()->IgnoreImpCasts()) &&
"Base should be the current struct!"); "Base should be the current struct!");
MatchingVD = ME->getMemberDecl(); MatchingVD = ME->getMemberDecl();
} }
// If we don't have information about the current list item, move on to // If we don't have information about the current list item, move on to
// the next one. // the next one.
auto InitAddrIt = CaptureDeviceAddrMap.find(MatchingVD); auto InitAddrIt = CaptureDeviceAddrMap.find(MatchingVD);
if (InitAddrIt == CaptureDeviceAddrMap.end()) if (InitAddrIt == CaptureDeviceAddrMap.end())
▲ Show 20 Lines • Show All 717 Lines • Show Last 20 Lines

clang/test/OpenMP/target_data_use_device_ptr_inheritance_codegen.cpp

  • This file was added.
// 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 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK1
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s -Wno-openmp-mapping
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK1
// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK2
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s -Wno-openmp-mapping
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK2
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
class A {
public:
double *ptr = nullptr;
virtual void foo() = 0;
};
class B : public A {
public:
virtual void foo() override;
};
void B::foo() {
#pragma omp target data use_device_ptr(A::ptr)
{
}
}
void foo() {
B b;
b.foo();
}
#endif
// CHECK-LABEL: define {{[^@]+}}@_ZN1B3fooEv
// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR0:[0-9]+]] align 2 {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8
// CHECK-NEXT: [[PTR4:%.*]] = alloca ptr, align 8
// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8
// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
// CHECK-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
// CHECK-NEXT: [[PTR3:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR3]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 8
// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 8
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
// CHECK-NEXT: store ptr null, ptr [[TMP3]], align 8
// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-NEXT: store ptr [[PTR2]], ptr [[TMP4]], align 8
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP5]], align 8
// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
// CHECK-NEXT: store ptr null, ptr [[TMP6]], align 8
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP7]], ptr [[TMP8]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
// CHECK-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP4]], align 8
// CHECK-NEXT: store ptr [[TMP9]], ptr [[PTR4]], align 8
// CHECK-NEXT: store ptr [[PTR4]], ptr [[TMP]], align 8
// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 2, ptr [[TMP10]], ptr [[TMP11]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
// CHECK-NEXT: ret void
// CHECK-LABEL: define {{[^@]+}}@_Z3foov
// CHECK-SAME: () #[[ATTR0]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[B:%.*]] = alloca [[CLASS_B:%.*]], align 8
// CHECK-NEXT: call void @_ZN1BC1Ev(ptr noundef nonnull align 8 dereferenceable(16) [[B]]) #[[ATTR1:[0-9]+]]
// CHECK-NEXT: call void @_ZN1B3fooEv(ptr noundef nonnull align 8 dereferenceable(16) [[B]])
// CHECK-NEXT: ret void
// CHECK-LABEL: define {{[^@]+}}@_ZN1BC1Ev
// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat align 2 {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: call void @_ZN1BC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
// CHECK-NEXT: ret void
// CHECK-LABEL: define {{[^@]+}}@_ZN1BC2Ev
// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: call void @_ZN1AC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
// CHECK-NEXT: store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1B, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
// CHECK-NEXT: ret void
// CHECK-LABEL: define {{[^@]+}}@_ZN1AC2Ev
// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1A, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
// CHECK-NEXT: [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
// CHECK-NEXT: store ptr null, ptr [[PTR]], align 8
// CHECK-NEXT: ret void
// CHECK1-LABEL: define {{[^@]+}}@_ZN1B3fooEv
// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR0:[0-9]+]] align 2 {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8
// CHECK1-NEXT: [[PTR4:%.*]] = alloca ptr, align 8
// CHECK1-NEXT: [[TMP:%.*]] = alloca ptr, align 8
// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
// CHECK1-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
// CHECK1-NEXT: [[PTR3:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR3]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 8
// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK1-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 8
// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
// CHECK1-NEXT: store ptr null, ptr [[TMP3]], align 8
// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK1-NEXT: store ptr [[PTR2]], ptr [[TMP4]], align 8
// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK1-NEXT: store ptr [[TMP0]], ptr [[TMP5]], align 8
// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
// CHECK1-NEXT: store ptr null, ptr [[TMP6]], align 8
// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK1-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP7]], ptr [[TMP8]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
// CHECK1-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP4]], align 8
// CHECK1-NEXT: store ptr [[TMP9]], ptr [[PTR4]], align 8
// CHECK1-NEXT: store ptr [[PTR4]], ptr [[TMP]], align 8
// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK1-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 2, ptr [[TMP10]], ptr [[TMP11]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
// CHECK1-NEXT: ret void
//
//
// CHECK1-LABEL: define {{[^@]+}}@_Z3foov
// CHECK1-SAME: () #[[ATTR0]] {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[B:%.*]] = alloca [[CLASS_B:%.*]], align 8
// CHECK1-NEXT: call void @_ZN1BC1Ev(ptr noundef nonnull align 8 dereferenceable(16) [[B]]) #[[ATTR1:[0-9]+]]
// CHECK1-NEXT: call void @_ZN1B3fooEv(ptr noundef nonnull align 8 dereferenceable(16) [[B]])
// CHECK1-NEXT: ret void
//
//
// CHECK1-LABEL: define {{[^@]+}}@_ZN1BC1Ev
// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat align 2 {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: call void @_ZN1BC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
// CHECK1-NEXT: ret void
//
//
// CHECK1-LABEL: define {{[^@]+}}@_ZN1BC2Ev
// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: call void @_ZN1AC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
// CHECK1-NEXT: store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1B, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
// CHECK1-NEXT: ret void
//
//
// CHECK1-LABEL: define {{[^@]+}}@_ZN1AC2Ev
// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK1-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK1-NEXT: store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1A, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
// CHECK1-NEXT: [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
// CHECK1-NEXT: store ptr null, ptr [[PTR]], align 8
// CHECK1-NEXT: ret void
//
//
// CHECK2-LABEL: define {{[^@]+}}@_ZN1B3fooEv
// CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR0:[0-9]+]] align 2 {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK2-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 4
// CHECK2-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 4
// CHECK2-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 4
// CHECK2-NEXT: [[PTR4:%.*]] = alloca ptr, align 4
// CHECK2-NEXT: [[TMP:%.*]] = alloca ptr, align 4
// CHECK2-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
// CHECK2-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
// CHECK2-NEXT: [[PTR3:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR3]], align 4
// CHECK2-NEXT: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK2-NEXT: store ptr [[THIS1]], ptr [[TMP1]], align 4
// CHECK2-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK2-NEXT: store ptr [[THIS1]], ptr [[TMP2]], align 4
// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
// CHECK2-NEXT: store ptr null, ptr [[TMP3]], align 4
// CHECK2-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK2-NEXT: store ptr [[PTR2]], ptr [[TMP4]], align 4
// CHECK2-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK2-NEXT: store ptr [[TMP0]], ptr [[TMP5]], align 4
// CHECK2-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1
// CHECK2-NEXT: store ptr null, ptr [[TMP6]], align 4
// CHECK2-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK2-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK2-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP7]], ptr [[TMP8]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
// CHECK2-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP4]], align 4
// CHECK2-NEXT: store ptr [[TMP9]], ptr [[PTR4]], align 4
// CHECK2-NEXT: store ptr [[PTR4]], ptr [[TMP]], align 4
// CHECK2-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK2-NEXT: [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK2-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 2, ptr [[TMP10]], ptr [[TMP11]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
// CHECK2-NEXT: ret void
//
//
// CHECK2-LABEL: define {{[^@]+}}@_Z3foov
// CHECK2-SAME: () #[[ATTR0]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[B:%.*]] = alloca [[CLASS_B:%.*]], align 4
// CHECK2-NEXT: call void @_ZN1BC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[B]]) #[[ATTR1:[0-9]+]]
// CHECK2-NEXT: call void @_ZN1B3fooEv(ptr noundef nonnull align 4 dereferenceable(8) [[B]])
// CHECK2-NEXT: ret void
//
//
// CHECK2-LABEL: define {{[^@]+}}@_ZN1BC1Ev
// CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat align 2 {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK2-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: call void @_ZN1BC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]]) #[[ATTR1]]
// CHECK2-NEXT: ret void
//
//
// CHECK2-LABEL: define {{[^@]+}}@_ZN1BC2Ev
// CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK2-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: call void @_ZN1AC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]]) #[[ATTR1]]
// CHECK2-NEXT: store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1B, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 4
// CHECK2-NEXT: ret void
//
//
// CHECK2-LABEL: define {{[^@]+}}@_ZN1AC2Ev
// CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK2-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
// CHECK2-NEXT: store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1A, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 4
// CHECK2-NEXT: [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
// CHECK2-NEXT: store ptr null, ptr [[PTR]], align 4
// CHECK2-NEXT: ret void
//