OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 2
Code Issues Packages Projects Releases Wiki Activity

[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
This commit is contained in:
Jennifer Yu 2022-08-11 11:23:57 -07:00
parent 4b33ea052a
commit 2ca27206f9
3 changed files with 63 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -9052,7 +9052,7 @@ public:
// 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 @@ public:
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 *>();

32
clang/test/OpenMP/target_is_device_ptr_codegen.cpp
View File

@ -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++;

31
openmp/libomptarget/test/mapping/is_device_ptr.cpp Normal file
View File

@ -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");
}