88 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 *arr;7 int y;8 int z;9};10 11int main() {12 S v;13 14#pragma omp target map(tofrom: v, v.x, v.z)15 {16 v.x++;17 v.y += 2;18 v.z += 3;19 }20 21#pragma omp target map(tofrom: v, v.x, v.arr[:1])22 {23 v.x++;24 v.y += 2;25 v.arr[0] += 2;26 v.z += 4;27 }28 29#pragma omp target map(tofrom: v, v.arr[:1])30 {31 v.x++;32 v.y += 2;33 v.arr[0] += 2;34 v.z += 4;35 }36 37 return 0;38}39 40// CHECK: [[CSTSZ0:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 4, i64 4]41// CHECK: [[CSTTY0:@.+]] = private {{.*}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]42 43// CHECK: [[CSTSZ1:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 4, i64 4]44// CHECK: [[CSTTY1:@.+]] = private {{.*}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]]]45 46// CHECK: [[CSTSZ2:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 24, i64 4]47// CHECK: [[CSTTY2:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]]]48 49// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])50// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 451// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 852// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [4 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 053 54// Fill two non-constant size elements here: the whole struct size, and the55// region covering v.arr and v.y.56 57// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 058// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 859// CHECK-DAG: [[ARRY:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 160// CHECK-DAG: store i64 %{{.+}}, ptr [[ARRY]], align 861 62// CHECK: call void63 64// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])65// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 466// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 867// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [4 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 068 69// Fill two non-constant size elements here: the whole struct size, and the70// region covering v.arr, v.y and v.z.71 72// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 073// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 874// CHECK-DAG: [[ARRYZ:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 175// CHECK-DAG: store i64 %{{.+}}, ptr [[ARRYZ]], align 876 77// CHECK: call void78 79// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])80// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 481// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 882// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [3 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 083 84// Fill one non-constant size element here: the whole struct size.85 86// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [3 x i64], ptr [[SIZES]], i32 0, i32 087// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 888