ABataev updated this revision to Diff 331623.
ABataev added a comment.
Rebase + fixes for references to complex data structures mapping
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D98812/new/
https://reviews.llvm.org/D98812
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_map_codegen_28.cpp
clang/test/OpenMP/target_map_codegen_35.cpp
openmp/libomptarget/test/mapping/data_member_ref.cpp
Index: openmp/libomptarget/test/mapping/data_member_ref.cpp
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/data_member_ref.cpp
@@ -0,0 +1,69 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <stdio.h>
+
+struct View {
+ int Data;
+};
+
+struct ViewPtr {
+ int *Data;
+};
+
+template <typename T> struct Foo {
+ Foo(T &V) : VRef(V) {}
+ T &VRef;
+};
+
+int main() {
+ View V;
+ V.Data = 123456;
+ Foo<View> Bar(V);
+ ViewPtr V1;
+ int Data = 123456;
+ V1.Data = &Data;
+ Foo<ViewPtr> Baz(V1);
+
+ // CHECK: Host 123456.
+ printf("Host %d.\n", Bar.VRef.Data);
+#pragma omp target map(Bar.VRef)
+ {
+ // CHECK: Device 123456.
+ printf("Device %d.\n", Bar.VRef.Data);
+ V.Data = 654321;
+ // CHECK: Device 654321.
+ printf("Device %d.\n", Bar.VRef.Data);
+ }
+ // CHECK: Host 654321 654321.
+ printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
+ V.Data = 123456;
+ // CHECK: Host 123456.
+ printf("Host %d.\n", Bar.VRef.Data);
+#pragma omp target map(Bar) map(Bar.VRef)
+ {
+ // CHECK: Device 123456.
+ printf("Device %d.\n", Bar.VRef.Data);
+ V.Data = 654321;
+ // CHECK: Device 654321.
+ printf("Device %d.\n", Bar.VRef.Data);
+ }
+ // CHECK: Host 654321 654321.
+ printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
+ // CHECK: Host 123456.
+ printf("Host %d.\n", *Baz.VRef.Data);
+#pragma omp target map(*Baz.VRef.Data)
+ {
+ // CHECK: Device 123456.
+ printf("Device %d.\n", *Baz.VRef.Data);
+ *V1.Data = 654321;
+ // CHECK: Device 654321.
+ printf("Device %d.\n", *Baz.VRef.Data);
+ }
+ // CHECK: Host 654321 654321 654321.
+ printf("Host %d %d %d.\n", *Baz.VRef.Data, *V1.Data, Data);
+ return 0;
+}
Index: clang/test/OpenMP/target_map_codegen_35.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/target_map_codegen_35.cpp
@@ -0,0 +1,182 @@
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-64
+// RUN: %clang_cc1 -DCK35 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-64
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-32
+// RUN: %clang_cc1 -DCK35 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-32
+
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK35 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK35 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
+#ifdef CK35
+
+class S {
+public:
+ S(double &b) : b(b) {}
+ int a;
+ double &b;
+ void foo();
+};
+
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | TO = 0x1000000000001
+// MEMBER_OF_1 | PTR_AND_OBJ | TO = 0x1000000000011
+// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]]
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | PTR_AND_OBJ | FROM = 0x1000000000012
+// CK35-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [2 x i64] [i64 [[#0x20]], i64 [[#0x1000000000012]]]
+
+void ref_map() {
+ double b;
+ S s(b);
+
+ // CK35-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null, i8** null)
+ // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)}
+
+ // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+ // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to %class.S**
+
+ // CK35-DAG: store %class.S* [[S_ADDR:%.+]], %class.S** [[BPC0]],
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC0]],
+ // CK35-DAG: store i64 [[S_SIZE:%.+]], i64* [[S0]],
+
+ // CK35-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK35-DAG: [[S_1_INTPTR]] = ptrtoint i8* [[S_1_VOID:%.+]] to i64
+ // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK35-DAG: [[S_1_VOID]] = bitcast %class.S* [[S_1:%.+]] to i8*
+ // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK35-DAG: [[S_1]] = getelementptr %class.S, %class.S* [[S_ADDR]], i32 1
+
+ // pass MEMBER_OF_1 | TO {&s, &s, ((void*)(&s.a+1)-(void*)&s)} to copy the data of s.a.
+
+ // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+
+ // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+ // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to %class.S**
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC1]],
+ // CK35-DAG: store i64 [[A_SIZE:%.+]], i64* [[S1]],
+
+ // CK35-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[B_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK35-DAG: [[B_BEGIN_INTPTR]] = ptrtoint i8* [[B_BEGIN_VOID:%.+]] to i64
+ // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK35-DAG: [[B_BEGIN_VOID]] = bitcast double** [[B_ADDR:%.+]] to i8*
+ // CK35-DAG: [[B_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ // pass MEMBER_OF_1 | TO {&s, &s.b+1, ((void*)(&s+1)-(void*)(&s.b+1))} to copy the data of remainder of s.
+
+ // CK35-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK35-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK35-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+
+ // CK35-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to %class.S**
+ // CK35-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double***
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC2]],
+ // CK35-DAG: store double** [[B_END:%.+]], double*** [[PC2]],
+ // CK35-DAG: store i64 [[REM_SIZE:%.+]], i64* [[S2]],
+
+ // CK35-DAG: [[B_END]] = getelementptr double*, double** [[B_ADDR]], i{{.+}} 1
+
+ // CK35-DAG: [[REM_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[B_END_INTPTR:%.+]]
+ // CK35-DAG: [[B_END_INTPTR]] = ptrtoint i8* [[B_END_VOID:%.+]] to i64
+ // CK35-DAG: [[S_END_INTPTR]] = ptrtoint i8* [[S_END_VOID:%.+]] to i64
+ // CK35-DAG: [[B_END_VOID]] = bitcast double** [[B_END]] to i8*
+ // CK35-DAG: [[S_END_VOID]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{.+}} 1
+ // CK35-64-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i64 15
+ // CK35-32-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i32 7
+ // CK35-DAG: [[S_VOIDPTR]] = bitcast %class.S* [[S_ADDR]] to i8*
+
+ // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b.
+
+ // CK35-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+ // CK35-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+ // CK35-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+
+ // CK35-DAG: [[BPC3:%.+]] = bitcast i8** [[BP3]] to %class.S**
+ // CK35-DAG: [[PC3:%.+]] = bitcast i8** [[P3]] to double**
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC3]],
+ // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC3]],
+ // CK35-DAG: store i64 8, i64* [[S3]],
+
+ // CK35-DAG: [[B_ADDR]] = load double*, double** [[B_REF:%.+]],
+ // CK35-DAG: [[B_REF]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ #pragma omp target map(to: s, s.b)
+ s.foo();
+
+ // CK35 : call void
+
+ // CK35-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null, i8** null)
+ // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // pass TARGET_PARAM {&s, &s.b, ((void*)(&s.b+1)-(void*)&s.b)}
+
+ // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+ // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double***
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC0]],
+ // CK35-DAG: store double** [[SB_ADDR:%.+]], double*** [[PC0]],
+ // CK35-DAG: store i64 [[B_SIZE:%.+]], i64* [[S0]],
+
+ // CK35-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[SB_1_INTPTR:%.+]], [[SB_INTPTR:%.+]]
+ // CK35-DAG: [[SB_1_INTPTR]] = ptrtoint i8* [[SB_1_VOID:%.+]] to i64
+ // CK35-DAG: [[SB_INTPTR]] = ptrtoint i8* [[SB_VOID:%.+]] to i64
+ // CK35-DAG: [[SB_1_VOID]] = bitcast double** [[SB_1:%.+]] to i8*
+ // CK35-DAG: [[SB_VOID]] = bitcast double** [[SB_ADDR:%.+]] to i8*
+ // CK35-DAG: [[SB_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+ // CK35-DAG: [[SB_1]] = getelementptr double*, double** [[SB_ADDR]], i{{.+}} 1
+
+ // pass MEMBER_OF_1 | PTR_AND_OBJ | FROM {&s, &s.b, 8|4} to copy the data of s.c.
+
+ // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+
+ // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+ // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double**
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+ // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC1]],
+ // CK35-DAG: store i64 8, i64* [[S1]],
+
+ // CK35-DAG: [[B_ADDR]] = load double*, double** [[SB_ADDR]],
+
+ #pragma omp target map(from: s.b)
+ s.foo();
+}
+
+#endif // CK35
+#endif
Index: clang/test/OpenMP/target_map_codegen_28.cpp
===================================================================
--- clang/test/OpenMP/target_map_codegen_28.cpp
+++ clang/test/OpenMP/target_map_codegen_28.cpp
@@ -70,7 +70,7 @@
// CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
- // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+ // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]***
// CK29-DAG: store [[SSB]]* [[VAR0:%.+]], [[SSB]]** [[CBP0]]
// CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
// CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
@@ -116,11 +116,10 @@
// CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
- // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+ // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]****
// CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]]
- // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
+ // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]]
// CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
- // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]],
// CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1
// CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
@@ -128,9 +127,10 @@
// CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
// CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
// CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
- // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
+ // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CBP1]]
// CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
// CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
+ // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000]],
// CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 0
// CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
@@ -161,11 +161,10 @@
// CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
- // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+ // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]****
// CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]]
- // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
+ // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]]
// CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
- // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]],
// CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1
// CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
@@ -173,9 +172,10 @@
// CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
// CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
// CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
- // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
+ // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CBP1]]
// CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
// CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
+ // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000]],
// CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]],
// CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7429,6 +7429,7 @@
// S1 s;
// double *p;
// struct S2 *ps;
+ // int &ref;
// }
// S2 s;
// S2 *ps;
@@ -7472,6 +7473,14 @@
// optimizes this entry out, same in the examples below)
// (***) map the pointee (map: to)
//
+ // map(to: s.ref)
+ // &s, &(s.ref), sizeof(int*), TARGET_PARAM (*)
+ // &s, &(s.ref), sizeof(int), MEMBER_OF(1) | PTR_AND_OBJ | TO (***)
+ // (*) alloc space for struct members, only this is a target parameter
+ // (**) map the pointer (nothing to be mapped in this example) (the compiler
+ // optimizes this entry out, same in the examples below)
+ // (***) map the pointee (map: to)
+ //
// map(s.ps)
// &s, &(s.ps), sizeof(S2*), TARGET_PARAM | TO | FROM
//
@@ -7669,6 +7678,7 @@
uint64_t DimSize = 1;
bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
+ bool IsPrevMemberReference = false;
for (; I != CE; ++I) {
// If the current component is member of a struct (parent struct) mark it.
@@ -7726,12 +7736,16 @@
.getCanonicalType()
->isAnyPointerType()) ||
I->getAssociatedExpression()->getType()->isAnyPointerType();
+ bool IsMemberReference = isa<MemberExpr>(I->getAssociatedExpression()) &&
+ MapDecl &&
+ MapDecl->getType()->isLValueReferenceType();
bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous;
if (OASE)
++DimSize;
- if (Next == CE || IsNonDerefPointer || IsFinalArraySection) {
+ if (Next == CE || IsMemberReference || IsNonDerefPointer ||
+ IsFinalArraySection) {
// If this is not the last component, we expect the pointer to be
// associated with an array expression or member expression.
assert((Next == CE ||
@@ -7744,22 +7758,39 @@
"Unexpected expression");
Address LB = Address::invalid();
+ Address LowestElem = Address::invalid();
if (OAShE) {
- LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
- CGF.getContext().getTypeAlignInChars(
- OAShE->getBase()->getType()));
- } else {
- LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
+ LowestElem = LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
+ CGF.getContext().getTypeAlignInChars(
+ OAShE->getBase()->getType()));
+ } else if (IsMemberReference) {
+ Address Base = BP;
+ QualType BaseType = EncounteredME->getBase()->getType();
+ if (EncounteredME->isArrow()) {
+ if (!IsExpressionFirstInfo || FirstPointerInComplexData)
+ Base =
+ CGF.EmitLoadOfPointer(Base, BaseType->castAs<PointerType>());
+ BaseType = BaseType->getPointeeType();
+ }
+ LValue BaseLVal = CGF.MakeAddrLValue(Base, BaseType);
+ LowestElem = CGF.EmitLValueForFieldInitialization(
+ BaseLVal, cast<FieldDecl>(MapDecl))
+ .getAddress(CGF);
+ LB = CGF.EmitLoadOfReferenceLValue(LowestElem, MapDecl->getType())
.getAddress(CGF);
+ } else {
+ LowestElem = LB =
+ CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
+ .getAddress(CGF);
}
// If this component is a pointer inside the base struct then we don't
// need to create any entry for it - it will be combined with the object
// it is pointing to into a single PTR_AND_OBJ entry.
bool IsMemberPointerOrAddr =
- (IsPointer || ForDeviceAddr) && EncounteredME &&
- (dyn_cast<MemberExpr>(I->getAssociatedExpression()) ==
- EncounteredME);
+ EncounteredME && (((IsPointer || ForDeviceAddr) &&
+ I->getAssociatedExpression() == EncounteredME) ||
+ (IsPrevMemberReference && !IsPointer));
if (!OverlappedElements.empty() && Next == CE) {
// Handle base element with the info for overlapped elements.
assert(!PartialStruct.Base.isValid() && "The base element is set.");
@@ -7767,11 +7798,11 @@
"Unexpected base element with the pointer type.");
// Mark the whole struct as the struct that requires allocation on the
// device.
- PartialStruct.LowestElem = {0, LB};
+ PartialStruct.LowestElem = {0, LowestElem};
CharUnits TypeSize = CGF.getContext().getTypeSizeInChars(
I->getAssociatedExpression()->getType());
Address HB = CGF.Builder.CreateConstGEP(
- CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LB,
+ CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LowestElem,
CGF.VoidPtrTy),
TypeSize.getQuantity() - 1);
PartialStruct.HighestElem = {
@@ -7797,10 +7828,28 @@
Address ComponentLB = Address::invalid();
for (const OMPClauseMappableExprCommon::MappableComponent &MC :
Component) {
- if (MC.getAssociatedDeclaration()) {
- ComponentLB =
- CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())
- .getAddress(CGF);
+ if (const ValueDecl *VD = MC.getAssociatedDeclaration()) {
+ const auto *FD = dyn_cast<FieldDecl>(VD);
+ if (FD && FD->getType()->isLValueReferenceType()) {
+ Address Base = BP;
+ const auto *ME =
+ cast<MemberExpr>(MC.getAssociatedExpression());
+ QualType BaseType = ME->getBase()->getType();
+ if (ME->isArrow()) {
+ if (!IsExpressionFirstInfo || FirstPointerInComplexData)
+ Base = CGF.EmitLoadOfPointer(
+ Base, BaseType->castAs<PointerType>());
+ BaseType = BaseType->getPointeeType();
+ }
+ LValue BaseLVal = CGF.MakeAddrLValue(Base, BaseType);
+ ComponentLB =
+ CGF.EmitLValueForFieldInitialization(BaseLVal, FD)
+ .getAddress(CGF);
+ } else {
+ ComponentLB =
+ CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())
+ .getAddress(CGF);
+ }
Size = CGF.Builder.CreatePtrDiff(
CGF.EmitCastToVoidPtr(ComponentLB.getPointer()),
CGF.EmitCastToVoidPtr(LB.getPointer()));
@@ -7856,13 +7905,13 @@
OpenMPOffloadMappingFlags Flags = getMapTypeBits(
MapType, MapModifiers, MotionModifiers, IsImplicit,
!IsExpressionFirstInfo || RequiresReference ||
- FirstPointerInComplexData,
+ FirstPointerInComplexData || IsMemberReference,
IsCaptureFirstInfo && !RequiresReference, IsNonContiguous);
- if (!IsExpressionFirstInfo) {
+ if (!IsExpressionFirstInfo || IsMemberReference) {
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
// then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
- if (IsPointer)
+ if (IsPointer || (IsMemberReference && Next != CE))
Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
OMP_MAP_DELETE | OMP_MAP_CLOSE);
@@ -7888,21 +7937,21 @@
// Update info about the lowest and highest elements for this struct
if (!PartialStruct.Base.isValid()) {
- PartialStruct.LowestElem = {FieldIndex, LB};
+ PartialStruct.LowestElem = {FieldIndex, LowestElem};
if (IsFinalArraySection) {
Address HB =
CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false)
.getAddress(CGF);
PartialStruct.HighestElem = {FieldIndex, HB};
} else {
- PartialStruct.HighestElem = {FieldIndex, LB};
+ PartialStruct.HighestElem = {FieldIndex, LowestElem};
}
PartialStruct.Base = BP;
PartialStruct.LB = BP;
} else if (FieldIndex < PartialStruct.LowestElem.first) {
- PartialStruct.LowestElem = {FieldIndex, LB};
+ PartialStruct.LowestElem = {FieldIndex, LowestElem};
} else if (FieldIndex > PartialStruct.HighestElem.first) {
- PartialStruct.HighestElem = {FieldIndex, LB};
+ PartialStruct.HighestElem = {FieldIndex, LowestElem};
}
}
@@ -7916,11 +7965,12 @@
// The pointer becomes the base for the next element.
if (Next != CE)
- BP = LB;
+ LowestElem = BP = LB;
IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
FirstPointerInComplexData = false;
+ IsPrevMemberReference = IsMemberReference;
} else if (FirstPointerInComplexData) {
QualType Ty = Components.rbegin()
->getAssociatedDeclaration()
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
