diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/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, bool, const ValueDecl *, const Expr *>; SmallVector 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() && "Expect a executable directive"); const auto *CurExecDir = CurDir.get(); diff --git a/clang/test/OpenMP/target_is_device_ptr_codegen.cpp b/clang/test/OpenMP/target_is_device_ptr_codegen.cpp --- a/clang/test/OpenMP/target_is_device_ptr_codegen.cpp +++ b/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++; @@ -268,15 +269,21 @@ // 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 @@ // 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 --- /dev/null +++ b/openmp/libomptarget/test/mapping/is_device_ptr.cpp @@ -0,0 +1,31 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include +#include + +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"); +}