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.

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

[OpenMP] Add OpenMP extension API to dump mapping tables #85381

Merged
merged 3 commits into from
Mar 18, 2024

Conversation

nicebert
Copy link
Contributor

This adds an API call ompx_dump_mapping_tables.
This allows users to debug the mapping tables and can be especially useful for unified shared memory applications to check if the code behaves in the way it should. The implementation reuses code already present to dump mapping tables (in a debug setting).

Copy link

Thank you for submitting a Pull Request (PR) to the LLVM Project!

This PR will be automatically labeled and the relevant teams will be
notified.

If you wish to, you can add reviewers by using the "Reviewers" section on this page.

If this is not working for you, it is probably because you do not have write
permissions for the repository. In which case you can instead tag reviewers by
name in a comment by using @ followed by their GitHub username.

If you have received no comments on your PR for a week, you can request a review
by "ping"ing the PR by adding a comment “Ping”. The common courtesy "ping" rate
is once a week. Please remember that you are asking for valuable time from other developers.

If you have further questions, they may be answered by the LLVM GitHub User Guide.

You can also ask questions in a comment on this PR, on the LLVM Discord or on the forums.

@llvmbot llvmbot added flang:openmp clang:openmp OpenMP related changes to Clang openmp:libomp OpenMP host runtime openmp:libomptarget OpenMP offload runtime labels Mar 15, 2024
@llvmbot
Copy link
Collaborator

llvmbot commented Mar 15, 2024

@llvm/pr-subscribers-flang-openmp

Author: None (nicebert)

Changes

This adds an API call ompx_dump_mapping_tables.
This allows users to debug the mapping tables and can be especially useful for unified shared memory applications to check if the code behaves in the way it should. The implementation reuses code already present to dump mapping tables (in a debug setting).


Full diff: https://github.com/llvm/llvm-project/pull/85381.diff

9 Files Affected:

  • (modified) llvm/include/llvm/Frontend/OpenMP/OMPKinds.def (+2)
  • (modified) openmp/libomptarget/include/OpenMP/Mapping.h (+1-1)
  • (modified) openmp/libomptarget/include/Shared/Debug.h (+17-5)
  • (modified) openmp/libomptarget/include/omptarget.h (+1)
  • (modified) openmp/libomptarget/src/OpenMP/API.cpp (+8)
  • (modified) openmp/libomptarget/src/OpenMP/Mapping.cpp (+7-5)
  • (modified) openmp/libomptarget/src/exports (+1)
  • (added) openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp (+37)
  • (modified) openmp/runtime/src/include/omp.h.var (+2)
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index d22d2a8e948b00..82ee551916a897 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -458,6 +458,8 @@ __OMP_RTL(__tgt_push_mapper_component, false, Void, VoidPtr, VoidPtr, VoidPtr,
 __OMP_RTL(__kmpc_task_allow_completion_event, false, VoidPtr, IdentPtr,
           /* Int */ Int32, /* kmp_task_t */ VoidPtr)
 
+__OMP_RTL(ompx_dump_mapping_tables, false, Void, )
+
 /// OpenMP Device runtime functions
 __OMP_RTL(__kmpc_target_init, false, Int32, KernelEnvironmentPtr, KernelLaunchEnvironmentPtr)
 __OMP_RTL(__kmpc_target_deinit, false, Void,)
diff --git a/openmp/libomptarget/include/OpenMP/Mapping.h b/openmp/libomptarget/include/OpenMP/Mapping.h
index 4bd676fc658a7d..e77f2bc2346544 100644
--- a/openmp/libomptarget/include/OpenMP/Mapping.h
+++ b/openmp/libomptarget/include/OpenMP/Mapping.h
@@ -424,7 +424,7 @@ typedef int (*TargetDataFuncPtrTy)(ident_t *, DeviceTy &, int32_t, void **,
                                    map_var_info_t *, void **, AsyncInfoTy &,
                                    bool);
 
-void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device);
+void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device, bool toStdOut = false);
 
 int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
                     void **ArgsBase, void **Args, int64_t *ArgSizes,
diff --git a/openmp/libomptarget/include/Shared/Debug.h b/openmp/libomptarget/include/Shared/Debug.h
index a39626d15386b0..ede5f845c27b8f 100644
--- a/openmp/libomptarget/include/Shared/Debug.h
+++ b/openmp/libomptarget/include/Shared/Debug.h
@@ -136,11 +136,13 @@ inline uint32_t getDebugLevel() {
   } while (0)
 
 /// Print a generic information string used if LIBOMPTARGET_INFO=1
-#define INFO_MESSAGE(_num, ...)                                                \
-  do {                                                                         \
-    fprintf(stderr, GETNAME(TARGET_NAME) " device %d info: ", (int)_num);      \
-    fprintf(stderr, __VA_ARGS__);                                              \
-  } while (0)
+#define INFO_MESSAGE(_num, ...) INFO_MESSAGE_TO(stderr, _num, __VA_ARGS__) 
+
+#define INFO_MESSAGE_TO(_stdDst, _num, ...)                                     \
+   do {                                                                         \
+      fprintf(_stdDst, GETNAME(TARGET_NAME) " device %d info: ", (int)_num);    \
+      fprintf(_stdDst, __VA_ARGS__);                                            \
+    } while (0)
 
 // Debugging messages
 #ifdef OMPTARGET_DEBUG
@@ -187,4 +189,14 @@ inline uint32_t getDebugLevel() {
     }                                                                          \
   } while (false)
 
+#define DUMP_INFO(always, _flags, _id, ...)                                    \
+  do {                                                                         \
+    if (always) {                                                              \
+      INFO_MESSAGE_TO(stdout, _id, __VA_ARGS__);                               \
+    } else {                                                                   \
+      INFO(_flags, _id, __VA_ARGS__);                                          \
+    }                                                                          \
+  } while (false)
+
+
 #endif // OMPTARGET_SHARED_DEBUG_H
diff --git a/openmp/libomptarget/include/omptarget.h b/openmp/libomptarget/include/omptarget.h
index 8e0ccf191839da..f01f7e90581309 100644
--- a/openmp/libomptarget/include/omptarget.h
+++ b/openmp/libomptarget/include/omptarget.h
@@ -273,6 +273,7 @@ struct __tgt_target_non_contig {
 extern "C" {
 #endif
 
+void ompx_dump_mapping_tables();
 int omp_get_num_devices(void);
 int omp_get_device_num(void);
 int omp_get_initial_device(void);
diff --git a/openmp/libomptarget/src/OpenMP/API.cpp b/openmp/libomptarget/src/OpenMP/API.cpp
index 85fb08c00a9a74..91104511e1bcb7 100644
--- a/openmp/libomptarget/src/OpenMP/API.cpp
+++ b/openmp/libomptarget/src/OpenMP/API.cpp
@@ -17,6 +17,7 @@
 
 #include "OpenMP/InternalTypes.h"
 #include "OpenMP/OMPT/Interface.h"
+#include "OpenMP/Mapping.h"
 #include "OpenMP/omp.h"
 #include "Shared/Profile.h"
 
@@ -27,6 +28,13 @@
 #include <cstring>
 #include <mutex>
 
+EXTERN void ompx_dump_mapping_tables() {
+  ident_t Loc = {0, 0, 0, 0, ";libomptarget;libomptarget;0;0;;"};
+  auto ExclusiveDevicesAccessor = PM->getExclusiveDevicesAccessor();
+  for (auto &Device : PM->devices(ExclusiveDevicesAccessor))
+    dumpTargetPointerMappings(&Loc, Device, true);
+}
+
 #ifdef OMPT_SUPPORT
 using namespace llvm::omp::target::ompt;
 #endif
diff --git a/openmp/libomptarget/src/OpenMP/Mapping.cpp b/openmp/libomptarget/src/OpenMP/Mapping.cpp
index 9c0b219b6f15f1..90a914800d68fc 100644
--- a/openmp/libomptarget/src/OpenMP/Mapping.cpp
+++ b/openmp/libomptarget/src/OpenMP/Mapping.cpp
@@ -16,23 +16,25 @@
 #include "device.h"
 
 /// Dump a table of all the host-target pointer pairs on failure
-void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device) {
+void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device, bool toStdOut) {
   MappingInfoTy::HDTTMapAccessorTy HDTTMap =
       Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor();
-  if (HDTTMap->empty())
+  if (HDTTMap->empty()) {
+    DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID, "OpenMP Host-Device pointer mappings table empty\n");
     return;
+  }
 
   SourceInfo Kernel(Loc);
-  INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
+  DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID,
        "OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n",
        Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn());
-  INFO(OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s %s\n",
+  DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s %s\n",
        "Host Ptr", "Target Ptr", "Size (B)", "DynRefCount", "HoldRefCount",
        "Declaration");
   for (const auto &It : *HDTTMap) {
     HostDataToTargetTy &HDTT = *It.HDTT;
     SourceInfo Info(HDTT.HstPtrName);
-    INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
+    DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID,
          DPxMOD " " DPxMOD " %-8" PRIuPTR " %-11s %-12s %s at %s:%d:%d\n",
          DPxPTR(HDTT.HstPtrBegin), DPxPTR(HDTT.TgtPtrBegin),
          HDTT.HstPtrEnd - HDTT.HstPtrBegin, HDTT.dynRefCountToStr().c_str(),
diff --git a/openmp/libomptarget/src/exports b/openmp/libomptarget/src/exports
index d5432a9eed380d..f95544ec8329c8 100644
--- a/openmp/libomptarget/src/exports
+++ b/openmp/libomptarget/src/exports
@@ -35,6 +35,7 @@ VERS1.0 {
     __tgt_push_mapper_component;
     __kmpc_push_target_tripcount;
     __kmpc_push_target_tripcount_mapper;
+    ompx_dump_mapping_tables;
     omp_get_mapped_ptr;
     omp_get_num_devices;
     omp_get_device_num;
diff --git a/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp b/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp
new file mode 100644
index 00000000000000..c24eaf58ac65ac
--- /dev/null
+++ b/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp
@@ -0,0 +1,37 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic 
+
+#include <cstdio>
+#include <omp.h>
+
+#define N 10
+
+int main() {
+  int *a = new int[N]; // mapped and released from device 0
+  int *b = new int[N]; // mapped to device 2
+
+  // CHECK: Mapping tables after target enter data:
+  // CHECK-NEXT: omptarget device 0 info: OpenMP Host-Device pointer mappings after block
+  // CHECK-NEXT: omptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
+  // CHECK-NEXT: omptarget device 0 info: {{(0x[0-9a-f]{16})}} {{(0x[0-9a-f]{16})}}
+  // CHECK-NEXT: omptarget device 1 info: OpenMP Host-Device pointer mappings table empty
+  // CHECK-NEXT: omptarget device 2 info: OpenMP Host-Device pointer mappings after block
+  // CHECK-NEXT: omptarget device 2 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
+  // CHECK-NEXT: omptarget device 2 info: {{(0x[0-9a-f]{16})}} {{(0x[0-9a-f]{16})}}
+  #pragma omp target enter data device(0) map(to:a[:N])
+  #pragma omp target enter data device(2) map(to:b[:N])
+  printf("Mapping tables after target enter data:\n");
+  ompx_dump_mapping_tables();
+
+  // CHECK: Mapping tables after target exit data for a:
+  // CHECK-NEXT: omptarget device 0 info: OpenMP Host-Device pointer mappings table empty
+  // CHECK-NEXT: omptarget device 1 info: OpenMP Host-Device pointer mappings table empty
+  // CHECK-NEXT: omptarget device 2 info: OpenMP Host-Device pointer mappings after block
+  // CHECK-NEXT: omptarget device 2 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
+  // CHECK-NEXT: omptarget device 2 info: {{(0x[0-9a-f]{16})}} {{(0x[0-9a-f]{16})}}
+  #pragma omp target exit data device(0) map(release:a[:N])
+  printf("\nMapping tables after target exit data for a:\n");
+  ompx_dump_mapping_tables();
+
+  return 0;
+}
+
diff --git a/openmp/runtime/src/include/omp.h.var b/openmp/runtime/src/include/omp.h.var
index a1488ae9d21c61..91d4f238f3393d 100644
--- a/openmp/runtime/src/include/omp.h.var
+++ b/openmp/runtime/src/include/omp.h.var
@@ -156,6 +156,8 @@
     /* OpenMP 5.1 interop */
     typedef intptr_t omp_intptr_t;
 
+    extern void __KAI_KMPC_CONVENTION ompx_dump_mapping_tables();
+
     /* 0..omp_get_num_interop_properties()-1 are reserved for implementation-defined properties */
     typedef enum omp_interop_property {
         omp_ipr_fr_id = -1,

Copy link

github-actions bot commented Mar 15, 2024

✅ With the latest revision this PR passed the C/C++ code formatter.

@nicebert nicebert force-pushed the AddAPIForDumpingTables branch 6 times, most recently from e26880f to 4b92353 Compare March 15, 2024 11:33
This adds an API call ompx_dump_mapping_tables.
This allows users to debug the mapping tables and can be especially useful for unified shared memory applications to check if the code behaves in the way it should.
The implementation reuses code already present to dump mapping tables (in a debug setting).
Add void to  parameter list.

Co-authored-by: Joseph Huber <huberjn@outlook.com>
openmp/libomptarget/include/omptarget.h Outdated Show resolved Hide resolved
Co-authored-by: Joseph Huber <huberjn@outlook.com>
@jhuber6
Copy link
Contributor

jhuber6 commented Mar 18, 2024

Thanks, I'll merge once the CI finishes.

@@ -16,28 +16,33 @@
#include "device.h"

/// Dump a table of all the host-target pointer pairs on failure
void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device) {
void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device,
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Oh yeah, this should be ToStdOut.

@jhuber6 jhuber6 merged commit 20f5bcf into llvm:main Mar 18, 2024
4 checks passed
Copy link

@nicebert Congratulations on having your first Pull Request (PR) merged into the LLVM Project!

Your changes will be combined with recent changes from other authors, then tested
by our build bots. If there is a problem with a build, you may recieve a report in an email or a comment on this PR.

Please check whether problems have been caused by your change specifically, as
the builds can include changes from many authors. It is not uncommon for your
change to be included in a build that fails due to someone else's changes, or
infrastructure issues.

How to do this, and the rest of the post-merge process, is covered in detail here.

If your change does cause a problem, it may be reverted, or you can revert it yourself.
This is a normal part of LLVM development. You can fix your changes and open a new PR to merge them again.

If you don't get any reports, no action is required from you. Your changes are working as expected, well done!

@jplehr
Copy link
Contributor

jplehr commented Mar 19, 2024

// CHECK-NEXT: omptarget device 2 info: {{(0x[0-9a-f]{16})}} {{(0x[0-9a-f]{16})}}
// clang-format on
#pragma omp target enter data device(0) map(to : a[ : N])
#pragma omp target enter data device(2) map(to : b[ : N])
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We have only 1 device in the buildbots, so this will fail.

chencha3 pushed a commit to chencha3/llvm-project that referenced this pull request Mar 23, 2024
This adds an API call ompx_dump_mapping_tables.
This allows users to debug the mapping tables and can be especially
useful for unified shared memory applications to check if the code
behaves in the way it should. The implementation reuses code already
present to dump mapping tables (in a debug setting).

---------

Co-authored-by: Joseph Huber <huberjn@outlook.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
clang:openmp OpenMP related changes to Clang flang:openmp openmp:libomp OpenMP host runtime openmp:libomptarget OpenMP offload runtime
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

5 participants