Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[OpenMP][Fix] Fix test initializations #74797

Closed
wants to merge 2 commits into from
Closed

[OpenMP][Fix] Fix test initializations #74797

wants to merge 2 commits into from
+313 −0

Conversation

doru1004
Copy link
Contributor

@doru1004 doru1004 commented Dec 8, 2023

Make sure arrays used in test are properly initialized.

@llvmbot llvmbot added clang Clang issues not falling into any other category clang:codegen clang:openmp OpenMP related changes to Clang openmp:libomptarget OpenMP offload runtime labels Dec 8, 2023
@llvmbot
Copy link
Collaborator

llvmbot commented Dec 8, 2023 •
edited

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-clang-codegen

Author: Gheorghe-Teodor Bercea (doru1004)

Changes

Make sure arrays used in test are properly initialized.

Patch is 20.59 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/74797.diff

4 Files Affected:

  • (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+22)
  • (added) clang/test/OpenMP/map_struct_ordering.cpp (+172)
  • (modified) openmp/libomptarget/test/offloading/back2back_distribute.c (+5)
  • (added) openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp (+114)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index d2be8141a3a4b3..84a6b36646897d 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7731,10 +7731,30 @@ class MappableExprsHandler {
               IsImplicit, Mapper, VarRef, ForDeviceAddr);
         };
 
+    // Sort all map clauses and make sure all the maps containing array
+    // sections are processed last.
+    llvm::SmallVector<const OMPMapClause *, 16> SortedMapClauses;
     for (const auto *Cl : Clauses) {
       const auto *C = dyn_cast<OMPMapClause>(Cl);
       if (!C)
         continue;
+      const auto *EI = C->getVarRefs().begin();
+      if (*EI && !isa<OMPArraySectionExpr>(*EI)) {
+        SortedMapClauses.emplace_back(C);
+      }
+    }
+    for (const auto *Cl : Clauses) {
+      const auto *C = dyn_cast<OMPMapClause>(Cl);
+      if (!C)
+        continue;
+      const auto *EI = C->getVarRefs().begin();
+      if (*EI && isa<OMPArraySectionExpr>(*EI)) {
+        SortedMapClauses.emplace_back(C);
+      }
+    }
+
+    // Iterate over all map clauses:
+    for (const OMPMapClause *C : SortedMapClauses) {
       MapKind Kind = Other;
       if (llvm::is_contained(C->getMapTypeModifiers(),
                              OMPC_MAP_MODIFIER_present))
@@ -7751,6 +7771,7 @@ class MappableExprsHandler {
         ++EI;
       }
     }
+
     for (const auto *Cl : Clauses) {
       const auto *C = dyn_cast<OMPToClause>(Cl);
       if (!C)
@@ -7767,6 +7788,7 @@ class MappableExprsHandler {
         ++EI;
       }
     }
+
     for (const auto *Cl : Clauses) {
       const auto *C = dyn_cast<OMPFromClause>(Cl);
       if (!C)
diff --git a/clang/test/OpenMP/map_struct_ordering.cpp b/clang/test/OpenMP/map_struct_ordering.cpp
new file mode 100644
index 00000000000000..035b39b5b12ab4
--- /dev/null
+++ b/clang/test/OpenMP/map_struct_ordering.cpp
@@ -0,0 +1,172 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" --prefix-filecheck-ir-name _ --version 4
+
+// RUN: %clang_cc1  -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+struct Descriptor {
+  int *datum;
+  long int x;
+  int xi;
+  long int arr[1][30];
+};
+
+int map_struct() {
+  Descriptor dat = Descriptor();
+  dat.xi = 3;
+  dat.arr[0][0] = 1;
+
+  #pragma omp target enter data map(to: dat.datum[:10]) map(to: dat)
+
+  #pragma omp target
+  {
+    dat.xi = 4;
+    dat.datum[dat.arr[0][0]] = dat.xi;
+  }
+
+  #pragma omp target exit data map(from: dat)
+
+  return dat.xi;
+}
+
+#endif
+// CHECK-LABEL: define dso_local noundef signext i32 @_Z10map_structv(
+// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DAT:%.*]] = alloca [[STRUCT_DESCRIPTOR:%.*]], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_PTRS5:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS6:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_PTRS8:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    call void @llvm.memset.p0.i64(ptr align 8 [[DAT]], i8 0, i64 264, i1 false)
+// CHECK-NEXT:    [[XI:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[DAT]], i32 0, i32 2
+// CHECK-NEXT:    store i32 3, ptr [[XI]], align 8
+// CHECK-NEXT:    [[ARR:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[DAT]], i32 0, i32 3
+// CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1 x [30 x i64]], ptr [[ARR]], i64 0, i64 0
+// CHECK-NEXT:    [[ARRAYIDX1:%.*]] = getelementptr inbounds [30 x i64], ptr [[ARRAYIDX]], i64 0, i64 0
+// CHECK-NEXT:    store i64 1, ptr [[ARRAYIDX1]], align 8
+// CHECK-NEXT:    [[DATUM:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[DAT]], i32 0, i32 0
+// CHECK-NEXT:    [[DATUM2:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[DAT]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DATUM2]], align 8
+// CHECK-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 0
+// CHECK-NEXT:    [[TMP1:%.*]] = getelementptr [[STRUCT_DESCRIPTOR]], ptr [[DAT]], i32 1
+// CHECK-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[TMP1]] to i64
+// CHECK-NEXT:    [[TMP3:%.*]] = ptrtoint ptr [[DAT]] to i64
+// CHECK-NEXT:    [[TMP4:%.*]] = sub i64 [[TMP2]], [[TMP3]]
+// CHECK-NEXT:    [[TMP5:%.*]] = sdiv exact i64 [[TMP4]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 24, i1 false)
+// CHECK-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP6]], align 8
+// CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP7]], align 8
+// CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// CHECK-NEXT:    store i64 [[TMP5]], ptr [[TMP8]], align 8
+// CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK-NEXT:    store ptr null, ptr [[TMP9]], align 8
+// CHECK-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP10]], align 8
+// CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP11]], align 8
+// CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// CHECK-NEXT:    store ptr null, ptr [[TMP12]], align 8
+// CHECK-NEXT:    [[TMP13:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
+// CHECK-NEXT:    store ptr [[DATUM]], ptr [[TMP13]], align 8
+// CHECK-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
+// CHECK-NEXT:    store ptr [[ARRAYIDX3]], ptr [[TMP14]], align 8
+// CHECK-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
+// CHECK-NEXT:    store ptr null, ptr [[TMP15]], align 8
+// CHECK-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP17:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP18:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// CHECK-NEXT:    call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 3, ptr [[TMP16]], ptr [[TMP17]], ptr [[TMP18]], ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK-NEXT:    [[TMP19:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP19]], align 8
+// CHECK-NEXT:    [[TMP20:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP20]], align 8
+// CHECK-NEXT:    [[TMP21:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS6]], i64 0, i64 0
+// CHECK-NEXT:    store ptr null, ptr [[TMP21]], align 8
+// CHECK-NEXT:    [[TMP22:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS4]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP23:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS5]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP24:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
+// CHECK-NEXT:    store i32 2, ptr [[TMP24]], align 4
+// CHECK-NEXT:    [[TMP25:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
+// CHECK-NEXT:    store i32 1, ptr [[TMP25]], align 4
+// CHECK-NEXT:    [[TMP26:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
+// CHECK-NEXT:    store ptr [[TMP22]], ptr [[TMP26]], align 8
+// CHECK-NEXT:    [[TMP27:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
+// CHECK-NEXT:    store ptr [[TMP23]], ptr [[TMP27]], align 8
+// CHECK-NEXT:    [[TMP28:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
+// CHECK-NEXT:    store ptr @.offload_sizes.1, ptr [[TMP28]], align 8
+// CHECK-NEXT:    [[TMP29:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
+// CHECK-NEXT:    store ptr @.offload_maptypes.2, ptr [[TMP29]], align 8
+// CHECK-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
+// CHECK-NEXT:    store ptr null, ptr [[TMP30]], align 8
+// CHECK-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
+// CHECK-NEXT:    store ptr null, ptr [[TMP31]], align 8
+// CHECK-NEXT:    [[TMP32:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
+// CHECK-NEXT:    store i64 0, ptr [[TMP32]], align 8
+// CHECK-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
+// CHECK-NEXT:    store i64 0, ptr [[TMP33]], align 8
+// CHECK-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
+// CHECK-NEXT:    store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP34]], align 4
+// CHECK-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
+// CHECK-NEXT:    store [3 x i32] zeroinitializer, ptr [[TMP35]], align 4
+// CHECK-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
+// CHECK-NEXT:    store i32 0, ptr [[TMP36]], align 4
+// CHECK-NEXT:    [[TMP37:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z10map_structv_l23.region_id, ptr [[KERNEL_ARGS]])
+// CHECK-NEXT:    [[TMP38:%.*]] = icmp ne i32 [[TMP37]], 0
+// CHECK-NEXT:    br i1 [[TMP38]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
+// CHECK:       omp_offload.failed:
+// CHECK-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z10map_structv_l23(ptr [[DAT]]) #[[ATTR3:[0-9]+]]
+// CHECK-NEXT:    br label [[OMP_OFFLOAD_CONT]]
+// CHECK:       omp_offload.cont:
+// CHECK-NEXT:    [[TMP39:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP39]], align 8
+// CHECK-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[TMP40]], align 8
+// CHECK-NEXT:    [[TMP41:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0
+// CHECK-NEXT:    store ptr null, ptr [[TMP41]], align 8
+// CHECK-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS8]], i32 0, i32 0
+// CHECK-NEXT:    call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP42]], ptr [[TMP43]], ptr @.offload_sizes.3, ptr @.offload_maptypes.4, ptr null, ptr null)
+// CHECK-NEXT:    [[XI10:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[DAT]], i32 0, i32 2
+// CHECK-NEXT:    [[TMP44:%.*]] = load i32, ptr [[XI10]], align 8
+// CHECK-NEXT:    ret i32 [[TMP44]]
+//
+//
+// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z10map_structv_l23(
+// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(264) [[DAT:%.*]]) #[[ATTR4:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DAT_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    store ptr [[DAT]], ptr [[DAT_ADDR]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DAT_ADDR]], align 8
+// CHECK-NEXT:    [[XI:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR:%.*]], ptr [[TMP0]], i32 0, i32 2
+// CHECK-NEXT:    store i32 4, ptr [[XI]], align 8
+// CHECK-NEXT:    [[XI1:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[TMP0]], i32 0, i32 2
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[XI1]], align 8
+// CHECK-NEXT:    [[DATUM:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[TMP0]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[DATUM]], align 8
+// CHECK-NEXT:    [[ARR:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[TMP0]], i32 0, i32 3
+// CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds [1 x [30 x i64]], ptr [[ARR]], i64 0, i64 0
+// CHECK-NEXT:    [[ARRAYIDX2:%.*]] = getelementptr inbounds [30 x i64], ptr [[ARRAYIDX]], i64 0, i64 0
+// CHECK-NEXT:    [[TMP3:%.*]] = load i64, ptr [[ARRAYIDX2]], align 8
+// CHECK-NEXT:    [[ARRAYIDX3:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[TMP3]]
+// CHECK-NEXT:    store i32 [[TMP1]], ptr [[ARRAYIDX3]], align 4
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define internal void @.omp_offloading.requires_reg(
+// CHECK-SAME: ) #[[ATTR5:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    call void @__tgt_register_requires(i64 1)
+// CHECK-NEXT:    ret void
+//
diff --git a/openmp/libomptarget/test/offloading/back2back_distribute.c b/openmp/libomptarget/test/offloading/back2back_distribute.c
index 750e39061974a3..17eb6127f021e6 100644
--- a/openmp/libomptarget/test/offloading/back2back_distribute.c
+++ b/openmp/libomptarget/test/offloading/back2back_distribute.c
@@ -22,6 +22,11 @@ int main(int argc, char *argv[]) {
   double * b = (double *) malloc(MAX_N * sizeof(double));
   double * c = (double *) malloc(MAX_N * sizeof(double));
 
+
+  for (int i = 0 ; i < MAX_N ; i++) {
+    d[i] = d_h[i] = a_h[i] = 0.0;
+  }
+
 #pragma omp target enter data map(to:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N])
 
   for (int n = 32 ; n < MAX_N ; n+=5000) {
diff --git a/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp b/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp
new file mode 100644
index 00000000000000..c7ce4bade8de9a
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp
@@ -0,0 +1,114 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+#include <stdio.h>
+#include <stdlib.h>
+
+struct Descriptor {
+  int *datum;
+  long int x;
+  int *more_datum;
+  int xi;
+  int val_datum, val_more_datum;
+  long int arr[1][30];
+  int val_arr;
+};
+
+int main() {
+  Descriptor dat = Descriptor();
+  dat.datum = (int *)malloc(sizeof(int) * 10);
+  dat.more_datum = (int *)malloc(sizeof(int) * 20);
+  dat.xi = 3;
+  dat.arr[0][0] = 1;
+
+  dat.datum[7] = 7;
+  dat.more_datum[17] = 17;
+
+  /// The struct is mapped with type 0x0 when the pointer fields are mapped.
+  /// The struct is also map explicitely by the user. The second mapping by
+  /// the user must not overwrite the mapping set up for the pointer fields
+  /// when mapping the struct happens after the mapping of the pointers.
+
+  // clang-format off
+  // CHECK: Libomptarget --> Entry  0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x0, Name=unknown
+  // CHECK: Libomptarget --> Entry  1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x1000000000001, Name=unknown
+  // CHECK: Libomptarget --> Entry  2: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=40, Type=0x1000000000011, Name=unknown
+  // CHECK: Libomptarget --> Entry  3: Base=[[MORE_DATUM_HST_PTR_BASE:0x.*]], Begin=[[MORE_DATUM_HST_PTR_BEGIN:0x.*]], Size=80, Type=0x1000000000011, Name=unknown
+  // clang-format on
+
+  /// The struct will be mapped in the same order as the above entries.
+
+  /// First argument is the struct itself and it will be mapped once.
+
+  // clang-format off
+  // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[DAT_HST_PTR_BASE]], Size=288)...
+  // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 288 with host pointer [[DAT_HST_PTR_BASE]].
+  // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[DAT_HST_PTR_BASE]], HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtAllocBegin=[[DAT_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1, HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Moving 288 bytes (hst:[[DAT_HST_PTR_BASE]]) -> (tgt:[[DAT_DEVICE_PTR_BASE]])
+  // clang-format on
+
+  /// Second argument is dat.datum:
+  // clang-format off
+  // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[DATUM_HST_PTR_BASE]], Size=40)...
+  // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 40 with host pointer [[DATUM_HST_PTR_BASE]].
+  // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[DATUM_HST_PTR_BASE]], HstPtrBegin=[[DATUM_HST_PTR_BASE]], TgtAllocBegin=[[DATUM_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DATUM_DEVICE_PTR_BASE]], Size=40, DynRefCount=1, HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Moving 40 bytes (hst:[[DATUM_HST_PTR_BASE]]) -> (tgt:[[DATUM_DEVICE_PTR_BASE]])
+  // clang-format on
+
+  /// Third argument is dat.more_datum:
+  // clang-format off
+  // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], Size=80)...
+  // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 80 with host pointer [[MORE_DATUM_HST_PTR_BEGIN]].
+  // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[MORE_DATUM_HST_PTR_BEGIN]], HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], TgtAllocBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN:0x.*]], TgtPtrBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN]], Size=80, DynRefCount=1, HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Moving 80 bytes (hst:[[MORE_DATUM_HST_PTR_BEGIN]]) -> (tgt:[[MORE_DATUM_DEVICE_PTR_BEGIN]])
+  // clang-format on
+
+#pragma omp target enter data map(to : dat.datum[ : 10])                       \
+    map(to : dat.more_datum[ : 20]) map(to : dat)
+
+  /// Checks induced by having a target region:
+  // clang-format off
+  // CHECK: Libomptarget --> Entry  0: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x223, Name=unknown
+  // CHECK: Libomptarget --> Mapping exists (implicit) with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=2 (incremented), HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Obtained target argument [[DAT_DEVICE_PTR_BASE]] from host pointer [[DAT_HST_PTR_BASE]]
+  // clang-format on
+
+#pragma omp target
+  {
+    dat.xi = 4;
+    dat.datum[7]++;
+    dat.more_datum[17]++;
+    dat.val_datum = dat.datum[7];
+    dat.val_more_datum = dat.more_datum[17];
+    dat.datum[dat.arr[0][0]] = dat.xi;
+    dat.val_arr = dat.datum[dat.arr[0][0]];
+  }
+
+  /// Post-target region checks:
+  // clang-format off
+  // CHECK: Libomptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1 (decremented), HoldRefCount=0
+  // clang-format on
+
+#pragma omp target exit data map(from : dat)
+
+  /// Target data end checks:
+  // clang-format off
+  // CHECK: Libomptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=0 (decremented, delay...
[truncated]

@doru1004 doru1004 closed this Dec 8, 2023
@doru1004 doru1004 deleted the fix-test-initializations branch December 8, 2023 02:32
@github-actions GitHub Actions
Copy link

github-actions bot commented Dec 8, 2023

⚠️ C/C++ code formatter, clang-format found issues in your code. ⚠️

You can test this locally with the following command: 
git-clang-format --diff 4162a9bca42a1152cdf4ae92ff7b90351c10f332 144d2ad42f5b8c9b89881d6238684cfd799a8270 -- clang/test/OpenMP/map_struct_ordering.cpp openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp clang/lib/CodeGen/CGOpenMPRuntime.cpp openmp/libomptarget/test/offloading/back2back_distribute.c
View the diff from clang-format here. 
diff --git a/openmp/libomptarget/test/offloading/back2back_distribute.c b/openmp/libomptarget/test/offloading/back2back_distribute.c
index 17eb6127f0..30cdd69776 100644
--- a/openmp/libomptarget/test/offloading/back2back_distribute.c
+++ b/openmp/libomptarget/test/offloading/back2back_distribute.c
@@ -22,8 +22,7 @@ int main(int argc, char *argv[]) {
   double * b = (double *) malloc(MAX_N * sizeof(double));
   double * c = (double *) malloc(MAX_N * sizeof(double));
 
-
-  for (int i = 0 ; i < MAX_N ; i++) {
+  for (int i = 0; i < MAX_N; i++) {
     d[i] = d_h[i] = a_h[i] = 0.0;
   }

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers
No reviews
Assignees
No one assigned
Labels
clang:codegen clang:openmp OpenMP related changes to Clang clang Clang issues not falling into any other category openmp:libomptarget OpenMP offload runtime
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.

None yet
2 participants
@doru1004 @llvmbot