jyu2 updated this revision to Diff 451959.
jyu2 added a comment.
Thanks Alexey. This to fix test format.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D129608/new/
https://reviews.llvm.org/D129608
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_is_device_ptr_codegen.cpp
openmp/libomptarget/test/mapping/is_device_ptr.cpp
Index: openmp/libomptarget/test/mapping/is_device_ptr.cpp
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/is_device_ptr.cpp
@@ -0,0 +1,32 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <assert.h>
+#include <iostream>
+#include <omp.h>
+
+struct view {
+ const int size = 10;
+ int *data_host;
+ int *data_device;
+ void foo() {
+ std::size_t bytes = size * sizeof(int);
+ const int host_id = omp_get_initial_device();
+ const int device_id = omp_get_default_device();
+ data_host = (int *)malloc(bytes);
+ data_device = (int *)omp_target_alloc(bytes, device_id);
+#pragma omp target teams distribute parallel for is_device_ptr(data_device)
+ for (int i = 0; i < size; ++i)
+ data_device[i] = i;
+ omp_target_memcpy(data_host, data_device, bytes, 0, 0, host_id, device_id);
+ for (int i = 0; i < size; ++i)
+ assert(data_host[i] == i);
+ }
+};
+
+int main() {
+ view a;
+ a.foo();
+ // CHECK: PASSED
+ printf("PASSED\n");
+}
+
Index: clang/test/OpenMP/target_is_device_ptr_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_is_device_ptr_codegen.cpp
+++ clang/test/OpenMP/target_is_device_ptr_codegen.cpp
@@ -252,12 +252,13 @@
// CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+// CK2-DAG: [[A:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1:%.+]], i32 0, i32 0
// CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
-// CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
-// CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
+// CK2-DAG: store [[ST]]* [[THIS1]], [[ST]]** [[CBP0]]
+// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double**
+// CK2-DAG: store double** [[A]], double*** [[CP0]]
#pragma omp target is_device_ptr(a)
{
a++;
@@ -271,12 +272,13 @@
// CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+// CK2-DAG: [[B:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
// CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
-// CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
-// CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
+// CK2-DAG: store %struct.ST* [[THIS1]], %struct.ST** [[CBP0]]
+// CH2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
+// CK2-DAG: store double*** [[B]], double**** [[CP0]]
#pragma omp target is_device_ptr(b)
{
b++;
@@ -290,12 +292,14 @@
// CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+// CK2-DAG: [[A8:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 0
+// CK2-DAG: [[B9:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
// CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
-// CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
-// CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
+// CK2-DAG: store %struct.ST* [[THIS1]], %struct.ST** [[CBP0]]
+// CH2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to to double***
+// CK2-DAG: store double** [[A8]], double*** [[TMP64:%.+]]
#pragma omp target is_device_ptr(a, b)
{
a++;
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9052,7 +9052,7 @@
// If this declaration appears in a is_device_ptr clause we just have to
// pass the pointer by value. If it is a reference to a declaration, we just
// pass its value.
- if (DevPointersMap.count(VD)) {
+ if (VD && DevPointersMap.count(VD)) {
CombinedInfo.Exprs.push_back(VD);
CombinedInfo.BasePointers.emplace_back(Arg, VD);
CombinedInfo.Pointers.push_back(Arg);
@@ -9071,6 +9071,14 @@
OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool,
const ValueDecl *, const Expr *>;
SmallVector<MapData, 4> DeclComponentLists;
+ // For member fields list in is_device_ptr, store it in
+ // DeclComponentLists for generating components info.
+ auto It = DevPointersMap.find(VD);
+ if (It != DevPointersMap.end())
+ for (const auto MCL : It->second)
+ DeclComponentLists.emplace_back(
+ MCL, OMPC_MAP_to, OMPC_MAP_MODIFIER_unknown, /*IsImpicit = */ true,
+ nullptr, nullptr);
assert(CurDir.is<const OMPExecutableDirective *>() &&
"Expect a executable directive");
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
