[clang] 2ca2720 - [OpenMP] Fix segmentation fault when data field is used in is_device_pt
Jennifer Yu via cfe-commits
cfe-commits at lists.llvm.org
Fri Aug 12 17:19:59 PDT 2022
Author: Jennifer Yu
Date: 2022-08-12T17:10:26-07:00
New Revision: 2ca27206f9739bffb50b999bdfb76fec1d7e62a8
URL: https://github.com/llvm/llvm-project/commit/2ca27206f9739bffb50b999bdfb76fec1d7e62a8
DIFF: https://github.com/llvm/llvm-project/commit/2ca27206f9739bffb50b999bdfb76fec1d7e62a8.diff
LOG: [OpenMP] Fix segmentation fault when data field is used in is_device_pt
Currently, the field just emit map info for this pointer variable. It is
failed at run time. For the fields, the PartialStruct is created and it
needs call to emitCombinedEntry which create the base that covers all
the pieces.
The change is to generate map info as regular fields.
Differential Revision: https://reviews.llvm.org/D129608
Added:
openmp/libomptarget/test/mapping/is_device_ptr.cpp
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_is_device_ptr_codegen.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 9fc7360e92472..a2d36e7df6f12 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9052,7 +9052,7 @@ class MappableExprsHandler {
// 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 @@ class MappableExprsHandler {
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 *>();
diff --git a/clang/test/OpenMP/target_is_device_ptr_codegen.cpp b/clang/test/OpenMP/target_is_device_ptr_codegen.cpp
index ee09c1b606133..6a5375a528a99 100644
--- a/clang/test/OpenMP/target_is_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_is_device_ptr_codegen.cpp
@@ -252,12 +252,13 @@ struct ST {
// 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++;
@@ -268,15 +269,21 @@ struct ST {
// CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
// CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
+// CK2-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CK2-DAG: store i64* [[SIZE:%.+]], i64** [[SARG]]
// CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+// CK2-DAG: [[S:%[^,]+]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK2-DAG: [[SIZE:%[^,]+]] = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i32 0, i32 0
+// CK2-DAG: store i64 [[S]], i64* [[SIZE]]
+// 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]]
+// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
+// CK2-DAG: store double*** [[B]], double**** [[CP0]]
#pragma omp target is_device_ptr(b)
{
b++;
@@ -287,15 +294,22 @@ struct ST {
// CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]]
// CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]]
+// CK2-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CK2-DAG: store i64* [[SIZE:%.+]], i64** [[SARG]]
// 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: [[S:%[^,]+]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK2-DAG: store i64 [[S]], i64* [[SIZE:%.+]]
+
// 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++;
diff --git a/openmp/libomptarget/test/mapping/is_device_ptr.cpp b/openmp/libomptarget/test/mapping/is_device_ptr.cpp
new file mode 100644
index 0000000000000..6433f822f9d53
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/is_device_ptr.cpp
@@ -0,0 +1,31 @@
+// 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");
+}
More information about the cfe-commits
mailing list