Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -9051,7 +9051,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); @@ -9070,6 +9070,16 @@ OpenMPMapClauseKind, ArrayRef, bool, const ValueDecl *, const Expr *>; SmallVector DeclComponentLists; + if (DevPointersMap.count(VD)) { + // 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, true, + nullptr, nullptr); + } assert(CurDir.is() && "Expect a executable directive"); const auto *CurExecDir = CurDir.get(); 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,12 @@ // 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: store double** %a, double*** [[TMP3:%.+]] #pragma omp target is_device_ptr(a) { a++; @@ -271,12 +271,12 @@ // 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]] +// CK2-DAG: store double*** [[B]], double**** [[TMP30:%.+]] #pragma omp target is_device_ptr(b) { b++; @@ -290,12 +290,13 @@ // 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]] +// CK2-DAG: store double** [[A8]], double*** [[TMP64:%.+]] #pragma omp target is_device_ptr(a, b) { a++; Index: openmp/libomptarget/test/mapping/is_device_ptr.cpp =================================================================== --- /dev/null +++ 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"); +}