154a128c65
Should be working downstream now This reverts commit 9b61ff210fdff752d5db55b128474e9990258488.
47 lines
1.7 KiB
C++
47 lines
1.7 KiB
C++
// 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;
|
|
};
|
|
|
|
struct T : public S {
|
|
int a;
|
|
int b;
|
|
int c;
|
|
};
|
|
|
|
int main() {
|
|
T v;
|
|
|
|
// This one should have no gap between v.z & v.a.
|
|
#pragma omp target map(tofrom: v, v.y, v.z, v.a)
|
|
{
|
|
v.y++;
|
|
v.z += 2;
|
|
v.a += 3;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4, i64 0]
|
|
// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], 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 [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
|
|
|
|
// Fill three non-constant size elements here: the whole struct size, the region
|
|
// covering v.x, and the region covering v.b and v.c.
|
|
|
|
// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0
|
|
// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
|
|
// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1
|
|
// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8
|
|
// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2
|
|
// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8
|