// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s
// expected-no-diagnostics

struct S {
  int x;
  int y;
  int z;
  int *p1;
  int *p2;
};

struct T : public S {
  int a;
  int b;
  int c;
};

int main() {
  T v;

// &v,       &v,                 sizeof(v),                         ALLOC | PARAM
// &v,       &v.z + sizeof(int), sizeof(v.p1 to v.p2),              TO | FROM | MEMBER_OF_1
// &v,       &v.c + sizeof(int), sizeof(T) - (&v.c + sizeof(int))), TO | FROM | MEMBER_OF_1
// &v,       &v.x,               sizeof(v.x),                       TO | FROM | MEMBER_OF_1
// &v,       &v.y,               sizeof(v.y),                       TO | FROM | MEMBER_OF_1
// &v,       &v.z,               sizeof(v.z),                       TO | FROM | MEMBER_OF_1
// &v,       &v.a,               sizeof(v.a),                       TO | FROM | MEMBER_OF_1
// &v,       &v.b,               sizeof(v.b),                       TO | FROM | MEMBER_OF_1
// &v,       &v.c,               sizeof(v.c),                       TO | FROM | MEMBER_OF_1
// &v.p1[0], &v.p1[0],           8 * sizeof(int),                   TO | FROM
// &v.p1,    &v.p1[0],           sizeof(v.p1),                      ATTACH
#pragma omp target map(tofrom: v, v.x, v.y, v.z, v.p1[:8], v.a, v.b, v.c)
  {
    v.x++;
    v.y += 2;
    v.z += 3;
    v.p1[0] += 4;
    v.a += 7;
    v.b += 5;
    v.c += 6;
  }

  return 0;
}

// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [12 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4, i64 4, i64 4, i64 4, i64 32, i64 8, i64 0]
// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [12 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x3]], i64 [[#0x8000]], i64 288]

// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [12 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0

// Check for filling of three non-constant size elements here: the whole struct
// size, the (padded) region covering p1 & p2, and the padding at the end of
// struct T.

// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [12 x i64], ptr [[SIZES]], i32 0, i32 0
// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
// CHECK-DAG: [[P1P2:%.+]] = getelementptr inbounds [12 x i64], ptr [[SIZES]], i32 0, i32 1
// CHECK-DAG: store i64 %{{.+}}, ptr [[P1P2]], align 8
// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [12 x i64], ptr [[SIZES]], i32 0, i32 2
// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8