ABataev created this revision.
ABataev added a reviewer: jdoerfert.
Herald added subscribers: guansong, yaxunl.
Herald added a project: clang.
ABataev requested review of this revision.
Herald added a subscriber: sstefan1.
The compiler should treat array subscript with base pointer as a first
pointer in complex data, it is used only for member expression with base
pointer.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D91660
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
Index: clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
@@ -0,0 +1,48 @@
+// Test host codegen.
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple
powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu
-emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple
powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu
-emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple
powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu
-std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple
i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - |
FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple
i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple
i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch
%t -verify %s -emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown
-fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple
powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu
-emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown
-fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify
%s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown
-fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown
-fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown
-fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s
-emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple
powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu
-emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple
powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu
-emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown
-fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify
%s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown
-fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck
--check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple
i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown
-fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s
-emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+#pragma omp declare target
+typedef struct {
+ int *arr;
+} MyObject;
+
+MyObject *objects;
+#pragma omp end declare target
+
+// CHECK-DAG: [[SIZES0:@.+]] = private unnamed_addr constant [1 x i64] [i64
{{8|4}}]
+// CHECK-DAG: [[MAPS0:@.+]] = private unnamed_addr constant [1 x i64] [i64 49]
+// CHECK-DAG: [[MAPS1:@.+]] = private unnamed_addr constant [2 x i64] [i64 32,
i64 281474976710673]
+// CHECK: @main
+int main(void) {
+// CHECK: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8**
%{{.+}}, i8** %{{.+}}, i64* getelementptr inbounds ([1 x i64], [1 x i64]*
[[SIZES0]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]*
[[MAPS0]], i32 0, i32 0), i8** null)
+#pragma omp target enter data map(to : objects [0:1])
+// CHECK: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8**
%{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* getelementptr inbounds ([2 x i64], [2
x i64]* [[MAPS1]], i32 0, i32 0), i8** null)
+#pragma omp target enter data map(to : objects[0].arr [0:1])
+
+ return 0;
+}
+#endif
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7870,6 +7870,10 @@
IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
FirstPointerInComplexData = false;
+ } else if (FirstPointerInComplexData) {
+ BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
+ .getAddress(CGF);
+ FirstPointerInComplexData = false;
}
}
Index: clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
@@ -0,0 +1,48 @@
+// Test host codegen.
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+#pragma omp declare target
+typedef struct {
+ int *arr;
+} MyObject;
+
+MyObject *objects;
+#pragma omp end declare target
+
+// CHECK-DAG: [[SIZES0:@.+]] = private unnamed_addr constant [1 x i64] [i64 {{8|4}}]
+// CHECK-DAG: [[MAPS0:@.+]] = private unnamed_addr constant [1 x i64] [i64 49]
+// CHECK-DAG: [[MAPS1:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710673]
+// CHECK: @main
+int main(void) {
+// CHECK: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[SIZES0]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPS0]], i32 0, i32 0), i8** null)
+#pragma omp target enter data map(to : objects [0:1])
+// CHECK: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** %{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPS1]], i32 0, i32 0), i8** null)
+#pragma omp target enter data map(to : objects[0].arr [0:1])
+
+ return 0;
+}
+#endif
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7870,6 +7870,10 @@
IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
FirstPointerInComplexData = false;
+ } else if (FirstPointerInComplexData) {
+ BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
+ .getAddress(CGF);
+ FirstPointerInComplexData = false;
}
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
