51 lines · cpp
1// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s2// expected-no-diagnostics3 4template<typename C>5struct S {6 C x;7 C y;8 char z; // Hidden padding after here...9};10 11template<typename C>12struct T : public S<C> {13 C a;14 C b;15 C c;16};17 18int main() {19 T<int> v;20 21#pragma omp target map(tofrom: v, v.y, v.z, v.a)22 {23 v.y++;24 v.z += 2;25 v.a += 3;26 }27 28 return 0;29}30 31// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 1, i64 4]32// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]33 34// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])35// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 436// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 837// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 038 39// Fill four non-constant size elements here: the whole struct size, the region40// covering v.x, the region covering padding after v.z and the region covering41// v.b and v.c.42 43// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 044// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 845// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 146// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 847// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 248// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 849// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 350// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 851