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 int *p1;9 int *p2;10};11 12struct T : public S {13 int a;14 int b;15 int c;16};17 18int main() {19 T v;20 21#pragma omp target map(tofrom: v, v.x, v.y, v.z, v.p1[:8], v.a, v.b, v.c)22 {23 v.x++;24 v.y += 2;25 v.z += 3;26 v.p1[0] += 4;27 v.a += 7;28 v.b += 5;29 v.c += 6;30 }31 32 return 0;33}34 35// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [10 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4, i64 32, i64 4, i64 4, i64 4]36// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [10 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]37 38// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])39// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 440// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 841// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [10 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 042 43// Check for filling of four non-constant size elements here: the whole struct44// size, the (padded) region covering p1 & p2, and the padding at the end of45// struct T.46 47// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 048// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 849// CHECK-DAG: [[P1P2:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 150// CHECK-DAG: store i64 %{{.+}}, ptr [[P1P2]], align 851// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 252// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 853