53 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 4struct S {5 int x;6 int y;7 int z;8};9 10struct M : public S {11 int mid;12};13 14struct T : public M {15 int a;16 int b;17 int c;18};19 20int main() {21 T v;22 23#pragma omp target map(tofrom: v, v.y, v.z, v.a)24 {25 v.y++;26 v.z += 2;27 v.a += 3;28 v.mid += 5;29 }30 31 return 0;32}33 34// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 4, i64 4]35// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]36 37// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])38// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 439// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 840// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 041 42// Fill four non-constant size elements here: the whole struct size, the region43// covering v.x, the region covering v.mid and the region covering v.b and v.c.44 45// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 046// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 847// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 148// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 849// CHECK-DAG: [[MID:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 250// CHECK-DAG: store i64 %{{.+}}, ptr [[MID]], align 851// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 352// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 853