clang/test/OpenMP/target_data_use_device_ptr_inheritance_codegen.cpp

   1 // 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 _
   2 // 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
   3 // 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
   4 // 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
   5 // 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
   6 // 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
   7 // 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
   8 // expected-no-diagnostics
   9 #ifndef HEADER
  10 #define HEADER
  11
  12 class A {
  13 public:
  14   double *ptr = nullptr;
  15   virtual void foo() = 0;
  16 };
  17
  18 class B : public A {
  19 public:
  20   virtual void foo() override;
  21 };
  22
  23 void B::foo() {
  24 #pragma omp target data use_device_ptr(A::ptr)
  25   {
  26   }
  27 }
  28
  29 void foo() {
  30   B b;
  31   b.foo();
  32 }
  33
  34 #endif
  35 // CHECK-LABEL: define {{[^@]+}}@_ZN1B3fooEv
  36 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR0:[0-9]+]] align 2 {
  37 // CHECK-NEXT:  entry:
  38 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
  39 // CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
  40 // CHECK-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
  41 // CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8
  42 // CHECK-NEXT:    [[PTR4:%.*]] = alloca ptr, align 8
  43 // CHECK-NEXT:    [[TMP:%.*]] = alloca ptr, align 8
  44 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
  45 // CHECK-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
  46 // CHECK-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
  47 // CHECK-NEXT:    [[PTR2:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
  48 // CHECK-NEXT:    [[PTR3:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
  49 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[PTR3]], align 8
  50 // CHECK-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
  51 // CHECK-NEXT:    store ptr [[THIS1]], ptr [[TMP1]], align 8
  52 // CHECK-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
  53 // CHECK-NEXT:    store ptr [[THIS1]], ptr [[TMP2]], align 8
  54 // CHECK-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
  55 // CHECK-NEXT:    store ptr null, ptr [[TMP3]], align 8
  56 // CHECK-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
  57 // CHECK-NEXT:    store ptr [[PTR2]], ptr [[TMP4]], align 8
  58 // CHECK-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
  59 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[TMP5]], align 8
  60 // CHECK-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
  61 // CHECK-NEXT:    store ptr null, ptr [[TMP6]], align 8
  62 // CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
  63 // CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
  64 // 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)
  65 // CHECK-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP4]], align 8
  66 // CHECK-NEXT:    store ptr [[TMP9]], ptr [[PTR4]], align 8
  67 // CHECK-NEXT:    store ptr [[PTR4]], ptr [[TMP]], align 8
  68 // CHECK-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
  69 // CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
  70 // 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)
  71 // CHECK-NEXT:    ret void
  72 // CHECK-LABEL: define {{[^@]+}}@_Z3foov
  73 // CHECK-SAME: () #[[ATTR0]] {
  74 // CHECK-NEXT:  entry:
  75 // CHECK-NEXT:    [[B:%.*]] = alloca [[CLASS_B:%.*]], align 8
  76 // CHECK-NEXT:    call void @_ZN1BC1Ev(ptr noundef nonnull align 8 dereferenceable(16) [[B]]) #[[ATTR1:[0-9]+]]
  77 // CHECK-NEXT:    call void @_ZN1B3fooEv(ptr noundef nonnull align 8 dereferenceable(16) [[B]])
  78 // CHECK-NEXT:    ret void
  79 // CHECK-LABEL: define {{[^@]+}}@_ZN1BC1Ev
  80 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat align 2 {
  81 // CHECK-NEXT:  entry:
  82 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
  83 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
  84 // CHECK-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
  85 // CHECK-NEXT:    call void @_ZN1BC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
  86 // CHECK-NEXT:    ret void
  87 // CHECK-LABEL: define {{[^@]+}}@_ZN1BC2Ev
  88 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
  89 // CHECK-NEXT:  entry:
  90 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
  91 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
  92 // CHECK-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
  93 // CHECK-NEXT:    call void @_ZN1AC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
  94 // CHECK-NEXT:    store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1B, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
  95 // CHECK-NEXT:    ret void
  96 // CHECK-LABEL: define {{[^@]+}}@_ZN1AC2Ev
  97 // CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
  98 // CHECK-NEXT:  entry:
  99 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 100 // CHECK-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 101 // CHECK-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 102 // CHECK-NEXT:    store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1A, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
 103 // CHECK-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
 104 // CHECK-NEXT:    store ptr null, ptr [[PTR]], align 8
 105 // CHECK-NEXT:    ret void
 106 // CHECK1-LABEL: define {{[^@]+}}@_ZN1B3fooEv
 107 // CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR0:[0-9]+]] {
 108 // CHECK1-NEXT:  entry:
 109 // CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 110 // CHECK1-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
 111 // CHECK1-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
 112 // CHECK1-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8
 113 // CHECK1-NEXT:    [[TMP0:%.*]] = alloca ptr, align 8
 114 // CHECK1-NEXT:    [[TMP:%.*]] = alloca ptr, align 8
 115 // CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 116 // CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 117 // CHECK1-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
 118 // CHECK1-NEXT:    [[PTR2:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
 119 // CHECK1-NEXT:    [[PTR3:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
 120 // CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[PTR3]], align 8
 121 // CHECK1-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 122 // CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP2]], align 8
 123 // CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 124 // CHECK1-NEXT:    store ptr [[THIS1]], ptr [[TMP3]], align 8
 125 // CHECK1-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
 126 // CHECK1-NEXT:    store ptr null, ptr [[TMP4]], align 8
 127 // CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
 128 // CHECK1-NEXT:    store ptr [[PTR2]], ptr [[TMP5]], align 8
 129 // CHECK1-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
 130 // CHECK1-NEXT:    store ptr [[TMP1]], ptr [[TMP6]], align 8
 131 // CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
 132 // CHECK1-NEXT:    store ptr null, ptr [[TMP7]], align 8
 133 // CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 134 // CHECK1-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 135 // CHECK1-NEXT:    call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
 136 // CHECK1-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[TMP5]], align 8
 137 // CHECK1-NEXT:    store ptr [[TMP10]], ptr [[TMP0]], align 8
 138 // CHECK1-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 8
 139 // CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 140 // CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 141 // CHECK1-NEXT:    call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 2, ptr [[TMP11]], ptr [[TMP12]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
 142 // CHECK1-NEXT:    ret void
 143 //
 144 //
 145 // CHECK1-LABEL: define {{[^@]+}}@_Z3foov
 146 // CHECK1-SAME: () #[[ATTR0]] {
 147 // CHECK1-NEXT:  entry:
 148 // CHECK1-NEXT:    [[B:%.*]] = alloca [[CLASS_B:%.*]], align 8
 149 // CHECK1-NEXT:    call void @_ZN1BC1Ev(ptr noundef nonnull align 8 dereferenceable(16) [[B]]) #[[ATTR1:[0-9]+]]
 150 // CHECK1-NEXT:    call void @_ZN1B3fooEv(ptr noundef nonnull align 8 dereferenceable(16) [[B]])
 151 // CHECK1-NEXT:    ret void
 152 //
 153 //
 154 // CHECK1-LABEL: define {{[^@]+}}@_ZN1BC1Ev
 155 // CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat {
 156 // CHECK1-NEXT:  entry:
 157 // CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 158 // CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 159 // CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 160 // CHECK1-NEXT:    call void @_ZN1BC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
 161 // CHECK1-NEXT:    ret void
 162 //
 163 //
 164 // CHECK1-LABEL: define {{[^@]+}}@_ZN1BC2Ev
 165 // CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
 166 // CHECK1-NEXT:  entry:
 167 // CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 168 // CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 169 // CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 170 // CHECK1-NEXT:    call void @_ZN1AC2Ev(ptr noundef nonnull align 8 dereferenceable(16) [[THIS1]]) #[[ATTR1]]
 171 // CHECK1-NEXT:    store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1B, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
 172 // CHECK1-NEXT:    ret void
 173 //
 174 //
 175 // CHECK1-LABEL: define {{[^@]+}}@_ZN1AC2Ev
 176 // CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
 177 // CHECK1-NEXT:  entry:
 178 // CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8
 179 // CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
 180 // CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
 181 // CHECK1-NEXT:    store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1A, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 8
 182 // CHECK1-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
 183 // CHECK1-NEXT:    store ptr null, ptr [[PTR]], align 8
 184 // CHECK1-NEXT:    ret void
 185 //
 186 //
 187 // CHECK2-LABEL: define {{[^@]+}}@_ZN1B3fooEv
 188 // CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR0:[0-9]+]] align 2 {
 189 // CHECK2-NEXT:  entry:
 190 // CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
 191 // CHECK2-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 4
 192 // CHECK2-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 4
 193 // CHECK2-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 4
 194 // CHECK2-NEXT:    [[TMP0:%.*]] = alloca ptr, align 4
 195 // CHECK2-NEXT:    [[TMP:%.*]] = alloca ptr, align 4
 196 // CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
 197 // CHECK2-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
 198 // CHECK2-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
 199 // CHECK2-NEXT:    [[PTR2:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
 200 // CHECK2-NEXT:    [[PTR3:%.*]] = getelementptr inbounds [[CLASS_A]], ptr [[THIS1]], i32 0, i32 1
 201 // CHECK2-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[PTR3]], align 4
 202 // CHECK2-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 203 // CHECK2-NEXT:    store ptr [[THIS1]], ptr [[TMP2]], align 4
 204 // CHECK2-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 205 // CHECK2-NEXT:    store ptr [[THIS1]], ptr [[TMP3]], align 4
 206 // CHECK2-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
 207 // CHECK2-NEXT:    store ptr null, ptr [[TMP4]], align 4
 208 // CHECK2-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
 209 // CHECK2-NEXT:    store ptr [[PTR2]], ptr [[TMP5]], align 4
 210 // CHECK2-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
 211 // CHECK2-NEXT:    store ptr [[TMP1]], ptr [[TMP6]], align 4
 212 // CHECK2-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1
 213 // CHECK2-NEXT:    store ptr null, ptr [[TMP7]], align 4
 214 // CHECK2-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 215 // CHECK2-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 216 // CHECK2-NEXT:    call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
 217 // CHECK2-NEXT:    [[TMP10:%.*]] = load ptr, ptr [[TMP5]], align 4
 218 // CHECK2-NEXT:    store ptr [[TMP10]], ptr [[TMP0]], align 4
 219 // CHECK2-NEXT:    store ptr [[TMP0]], ptr [[TMP]], align 4
 220 // CHECK2-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
 221 // CHECK2-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
 222 // CHECK2-NEXT:    call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 2, ptr [[TMP11]], ptr [[TMP12]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
 223 // CHECK2-NEXT:    ret void
 224 //
 225 //
 226 // CHECK2-LABEL: define {{[^@]+}}@_Z3foov
 227 // CHECK2-SAME: () #[[ATTR0]] {
 228 // CHECK2-NEXT:  entry:
 229 // CHECK2-NEXT:    [[B:%.*]] = alloca [[CLASS_B:%.*]], align 4
 230 // CHECK2-NEXT:    call void @_ZN1BC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[B]]) #[[ATTR1:[0-9]+]]
 231 // CHECK2-NEXT:    call void @_ZN1B3fooEv(ptr noundef nonnull align 4 dereferenceable(8) [[B]])
 232 // CHECK2-NEXT:    ret void
 233 //
 234 //
 235 // CHECK2-LABEL: define {{[^@]+}}@_ZN1BC1Ev
 236 // CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat align 2 {
 237 // CHECK2-NEXT:  entry:
 238 // CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
 239 // CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
 240 // CHECK2-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
 241 // CHECK2-NEXT:    call void @_ZN1BC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]]) #[[ATTR1]]
 242 // CHECK2-NEXT:    ret void
 243 //
 244 //
 245 // CHECK2-LABEL: define {{[^@]+}}@_ZN1BC2Ev
 246 // CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
 247 // CHECK2-NEXT:  entry:
 248 // CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
 249 // CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
 250 // CHECK2-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
 251 // CHECK2-NEXT:    call void @_ZN1AC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]]) #[[ATTR1]]
 252 // CHECK2-NEXT:    store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1B, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 4
 253 // CHECK2-NEXT:    ret void
 254 //
 255 //
 256 // CHECK2-LABEL: define {{[^@]+}}@_ZN1AC2Ev
 257 // CHECK2-SAME: (ptr noundef nonnull align 4 dereferenceable(8) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat align 2 {
 258 // CHECK2-NEXT:  entry:
 259 // CHECK2-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 4
 260 // CHECK2-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
 261 // CHECK2-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
 262 // CHECK2-NEXT:    store ptr getelementptr inbounds ({ [3 x ptr] }, ptr @_ZTV1A, i32 0, inrange i32 0, i32 2), ptr [[THIS1]], align 4
 263 // CHECK2-NEXT:    [[PTR:%.*]] = getelementptr inbounds [[CLASS_A:%.*]], ptr [[THIS1]], i32 0, i32 1
 264 // CHECK2-NEXT:    store ptr null, ptr [[PTR]], align 4
 265 // CHECK2-NEXT:    ret void
 266 //