[Openmp-commits] [openmp] [OpenMP] Fix ompx_dump_mapping_tables lit test (PR #85753)
via Openmp-commits
openmp-commits at lists.llvm.org
Tue Mar 19 01:33:38 PDT 2024
https://github.com/nicebert created https://github.com/llvm/llvm-project/pull/85753
Fixes ompx_dump_mapping_tables test by only using one device after breaking built bots
>From ed914ba2006be708f55f672f61ec695c779ed057 Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Fri, 15 Mar 2024 06:44:01 -0400
Subject: [PATCH 1/4] [OpenMP] Add OpenMP extension API to dump mapping tables
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).
---
openmp/libomptarget/include/OpenMP/Mapping.h | 3 +-
openmp/libomptarget/include/Shared/Debug.h | 17 ++++++--
openmp/libomptarget/include/omptarget.h | 1 +
openmp/libomptarget/src/OpenMP/API.cpp | 8 ++++
openmp/libomptarget/src/OpenMP/Mapping.cpp | 33 ++++++++-------
openmp/libomptarget/src/exports | 1 +
.../test/api/ompx_dump_mapping_tables.cpp | 40 +++++++++++++++++++
openmp/runtime/src/include/omp.h.var | 2 +
8 files changed, 87 insertions(+), 18 deletions(-)
create mode 100644 openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp
diff --git a/openmp/libomptarget/include/OpenMP/Mapping.h b/openmp/libomptarget/include/OpenMP/Mapping.h
index 4bd676fc658a7d..b9f5c165829314 100644
--- a/openmp/libomptarget/include/OpenMP/Mapping.h
+++ b/openmp/libomptarget/include/OpenMP/Mapping.h
@@ -424,7 +424,8 @@ 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..7c3db8dbf119f6 100644
--- a/openmp/libomptarget/include/Shared/Debug.h
+++ b/openmp/libomptarget/include/Shared/Debug.h
@@ -136,10 +136,12 @@ inline uint32_t getDebugLevel() {
} while (0)
/// Print a generic information string used if LIBOMPTARGET_INFO=1
-#define INFO_MESSAGE(_num, ...) \
+#define INFO_MESSAGE(_num, ...) INFO_MESSAGE_TO(stderr, _num, __VA_ARGS__)
+
+#define INFO_MESSAGE_TO(_stdDst, _num, ...) \
do { \
- fprintf(stderr, GETNAME(TARGET_NAME) " device %d info: ", (int)_num); \
- fprintf(stderr, __VA_ARGS__); \
+ fprintf(_stdDst, GETNAME(TARGET_NAME) " device %d info: ", (int)_num); \
+ fprintf(_stdDst, __VA_ARGS__); \
} while (0)
// Debugging messages
@@ -187,4 +189,13 @@ inline uint32_t getDebugLevel() {
} \
} while (false)
+#define DUMP_INFO(toStdOut, _flags, _id, ...) \
+ do { \
+ if (toStdOut) { \
+ 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..c85f9868e37c25 100644
--- a/openmp/libomptarget/src/OpenMP/API.cpp
+++ b/openmp/libomptarget/src/OpenMP/API.cpp
@@ -16,6 +16,7 @@
#include "rtl.h"
#include "OpenMP/InternalTypes.h"
+#include "OpenMP/Mapping.h"
#include "OpenMP/OMPT/Interface.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..d8ad5d8cacc4bb 100644
--- a/openmp/libomptarget/src/OpenMP/Mapping.cpp
+++ b/openmp/libomptarget/src/OpenMP/Mapping.cpp
@@ -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,
+ 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,
- "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",
- "Host Ptr", "Target Ptr", "Size (B)", "DynRefCount", "HoldRefCount",
- "Declaration");
+ 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());
+ 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,
- 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(),
- HDTT.holdRefCountToStr().c_str(), Info.getName(), Info.getFilename(),
- Info.getLine(), Info.getColumn());
+ 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(), HDTT.holdRefCountToStr().c_str(),
+ Info.getName(), Info.getFilename(), Info.getLine(),
+ Info.getColumn());
}
}
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..a57d0c8a6d2bf4
--- /dev/null
+++ b/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp
@@ -0,0 +1,40 @@
+// 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
+
+ // clang-format off
+ // 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})}}
+ // 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])
+ printf("Mapping tables after target enter data:\n");
+ ompx_dump_mapping_tables();
+
+ // clang-format off
+ // 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})}}
+ // clang-format on
+#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,
>From 78dde68aec8d683e6fc62389a30d957aa4bc695b Mon Sep 17 00:00:00 2001
From: nicebert <110385235+nicebert at users.noreply.github.com>
Date: Mon, 18 Mar 2024 16:48:33 +0100
Subject: [PATCH 2/4] Update openmp/runtime/src/include/omp.h.var
Add void to parameter list.
Co-authored-by: Joseph Huber <huberjn at outlook.com>
---
openmp/runtime/src/include/omp.h.var | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/openmp/runtime/src/include/omp.h.var b/openmp/runtime/src/include/omp.h.var
index 91d4f238f3393d..eb3ab7778606a3 100644
--- a/openmp/runtime/src/include/omp.h.var
+++ b/openmp/runtime/src/include/omp.h.var
@@ -156,7 +156,7 @@
/* OpenMP 5.1 interop */
typedef intptr_t omp_intptr_t;
- extern void __KAI_KMPC_CONVENTION ompx_dump_mapping_tables();
+ extern void __KAI_KMPC_CONVENTION ompx_dump_mapping_tables(void);
/* 0..omp_get_num_interop_properties()-1 are reserved for implementation-defined properties */
typedef enum omp_interop_property {
>From 678910806ad18cf114d29fe464f36151067f375c Mon Sep 17 00:00:00 2001
From: nicebert <110385235+nicebert at users.noreply.github.com>
Date: Mon, 18 Mar 2024 16:53:37 +0100
Subject: [PATCH 3/4] Update openmp/libomptarget/include/omptarget.h
Co-authored-by: Joseph Huber <huberjn at outlook.com>
---
openmp/libomptarget/include/omptarget.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/openmp/libomptarget/include/omptarget.h b/openmp/libomptarget/include/omptarget.h
index f01f7e90581309..323dee41630f2f 100644
--- a/openmp/libomptarget/include/omptarget.h
+++ b/openmp/libomptarget/include/omptarget.h
@@ -273,7 +273,7 @@ struct __tgt_target_non_contig {
extern "C" {
#endif
-void ompx_dump_mapping_tables();
+void ompx_dump_mapping_tables(void);
int omp_get_num_devices(void);
int omp_get_device_num(void);
int omp_get_initial_device(void);
>From 5730b7bfa3881a3c3dc75dee9b1381c07a37d7ee Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Tue, 19 Mar 2024 04:29:47 -0400
Subject: [PATCH 4/4] [OpenMP] Fix ompx_dump_mapping_tables lit test
Fixes ompx_dump_mapping_tables test by only using one device after breaking built bots.
---
.../test/api/ompx_dump_mapping_tables.cpp | 25 ++++++++-----------
1 file changed, 10 insertions(+), 15 deletions(-)
diff --git a/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp b/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp
index a57d0c8a6d2bf4..d9288204627007 100644
--- a/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp
+++ b/openmp/libomptarget/test/api/ompx_dump_mapping_tables.cpp
@@ -6,33 +6,28 @@
#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
+ int *a = new __int32_t[N]; // mapped and released from device 0
+ int *b = new __int32_t[2 * N]; // mapped to device 0
// clang-format off
// 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})}}
- // clang-format on
+ // CHECK-NEXT: omptarget device 0 info: {{(0x[0-9a-f]{16})}} {{(0x[0-9a-f]{16})}} 40
+ // CHECK-NEXT: omptarget device 0 info: {{(0x[0-9a-f]{16})}} {{(0x[0-9a-f]{16})}} 80
#pragma omp target enter data device(0) map(to : a[ : N])
-#pragma omp target enter data device(2) map(to : b[ : N])
+#pragma omp target enter data device(0) map(to : b[ : 2*N])
+ // clang-format on
printf("Mapping tables after target enter data:\n");
ompx_dump_mapping_tables();
// clang-format off
// 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})}}
- // clang-format on
+ // 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})}} 80
#pragma omp target exit data device(0) map(release : a[ : N])
+ // clang-format on
printf("\nMapping tables after target exit data for a:\n");
ompx_dump_mapping_tables();
More information about the Openmp-commits
mailing list