[llvm-branch-commits] [openmp] fe5d51a - [OpenMP] Add using bit flags to select Libomptarget Information
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Jan 4 09:09:58 PST 2021
Author: Joseph Huber
Date: 2021-01-04T12:03:15-05:00
New Revision: fe5d51a4897c26696fede55e120c912df60cd3f4
URL: https://github.com/llvm/llvm-project/commit/fe5d51a4897c26696fede55e120c912df60cd3f4
DIFF: https://github.com/llvm/llvm-project/commit/fe5d51a4897c26696fede55e120c912df60cd3f4.diff
LOG: [OpenMP] Add using bit flags to select Libomptarget Information
Summary:
This patch adds more fine-grained support over which information is output from the libomptarget runtime when run with the environment variable LIBOMPTARGET_INFO set. An extensible set of flags can be used to pick and choose which information the user is interested in.
Reviewers: jdoerfert JonChesterfield grokos
Differential Revision: https://reviews.llvm.org/D93727
Added:
Modified:
openmp/libomptarget/include/Debug.h
openmp/libomptarget/include/SourceInfo.h
openmp/libomptarget/plugins/cuda/src/rtl.cpp
openmp/libomptarget/src/device.cpp
openmp/libomptarget/src/interface.cpp
openmp/libomptarget/src/private.h
openmp/libomptarget/test/offloading/info.c
Removed:
################################################################################
diff --git a/openmp/libomptarget/include/Debug.h b/openmp/libomptarget/include/Debug.h
index 4f42794e1bca..de593ecf5c3e 100644
--- a/openmp/libomptarget/include/Debug.h
+++ b/openmp/libomptarget/include/Debug.h
@@ -37,24 +37,38 @@
#ifndef _OMPTARGET_DEBUG_H
#define _OMPTARGET_DEBUG_H
-static inline int getInfoLevel() {
- static int InfoLevel = -1;
- if (InfoLevel >= 0)
- return InfoLevel;
-
- if (char *EnvStr = getenv("LIBOMPTARGET_INFO"))
- InfoLevel = std::stoi(EnvStr);
+#include <mutex>
+
+/// 32-Bit field data attributes controlling information presented to the user.
+enum OpenMPInfoType : uint32_t {
+ // Print data arguments and attributes upon entering an OpenMP device kernel.
+ OMP_INFOTYPE_KERNEL_ARGS = 0x0001,
+ // Indicate when an address already exists in the device mapping table.
+ OMP_INFOTYPE_MAPPING_EXISTS = 0x0002,
+ // Dump the contents of the device pointer map at kernel exit or failure.
+ OMP_INFOTYPE_DUMP_TABLE = 0x0004,
+ // Print kernel information from target device plugins
+ OMP_INFOTYPE_PLUGIN_KERNEL = 0x0010,
+};
+
+static inline uint32_t getInfoLevel() {
+ static uint32_t InfoLevel = 0;
+ static std::once_flag Flag{};
+ std::call_once(Flag, []() {
+ if (char *EnvStr = getenv("LIBOMPTARGET_INFO"))
+ InfoLevel = std::stoi(EnvStr);
+ });
return InfoLevel;
}
-static inline int getDebugLevel() {
- static int DebugLevel = -1;
- if (DebugLevel >= 0)
- return DebugLevel;
-
- if (char *EnvStr = getenv("LIBOMPTARGET_DEBUG"))
- DebugLevel = std::stoi(EnvStr);
+static inline uint32_t getDebugLevel() {
+ static uint32_t DebugLevel = 0;
+ static std::once_flag Flag{};
+ std::call_once(Flag, []() {
+ if (char *EnvStr = getenv("LIBOMPTARGET_DEBUG"))
+ DebugLevel = std::stoi(EnvStr);
+ });
return DebugLevel;
}
@@ -107,7 +121,7 @@ static inline int getDebugLevel() {
/// Print a generic information string used if LIBOMPTARGET_INFO=1
#define INFO_MESSAGE(_num, ...) \
do { \
- fprintf(stderr, GETNAME(TARGET_NAME) " device %d info: ", _num); \
+ fprintf(stderr, GETNAME(TARGET_NAME) " device %d info: ", (int)_num); \
fprintf(stderr, __VA_ARGS__); \
} while (0)
diff --git a/openmp/libomptarget/include/SourceInfo.h b/openmp/libomptarget/include/SourceInfo.h
index 614f99e62afe..c659d916837b 100644
--- a/openmp/libomptarget/include/SourceInfo.h
+++ b/openmp/libomptarget/include/SourceInfo.h
@@ -54,6 +54,13 @@ class SourceInfo {
return std::string(reinterpret_cast<const char *>(name));
}
+ std::string initStr(const ident_t *loc) {
+ if (!loc)
+ return ";unknown;unknown;0;0;;";
+ else
+ return std::string(reinterpret_cast<const char *>(loc->psource));
+ }
+
/// Get n-th substring in an expression separated by ;.
std::string getSubstring(const int n) const {
std::size_t begin = sourceStr.find(';');
@@ -73,7 +80,7 @@ class SourceInfo {
public:
SourceInfo(const ident_t *loc)
- : sourceStr(initStr(loc->psource)), name(getSubstring(1)),
+ : sourceStr(initStr(loc)), name(getSubstring(1)),
filename(removePath(getSubstring(0))), line(std::stoi(getSubstring(2))),
column(std::stoi(getSubstring(3))) {}
diff --git a/openmp/libomptarget/plugins/cuda/src/rtl.cpp b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
index 0422bfbfe319..4fac6a76710e 100644
--- a/openmp/libomptarget/plugins/cuda/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
@@ -501,11 +501,12 @@ class DeviceRTLTy {
DeviceData[DeviceId].BlocksPerGrid = EnvTeamLimit;
}
- INFO(DeviceId,
- "Device supports up to %d CUDA blocks and %d threads with a "
- "warp size of %d\n",
- DeviceData[DeviceId].BlocksPerGrid,
- DeviceData[DeviceId].ThreadsPerBlock, DeviceData[DeviceId].WarpSize);
+ if (getDebugLevel() || (getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL))
+ INFO(DeviceId,
+ "Device supports up to %d CUDA blocks and %d threads with a "
+ "warp size of %d\n",
+ DeviceData[DeviceId].BlocksPerGrid,
+ DeviceData[DeviceId].ThreadsPerBlock, DeviceData[DeviceId].WarpSize);
// Set default number of teams
if (EnvNumTeams > 0) {
@@ -937,14 +938,15 @@ class DeviceRTLTy {
CudaBlocksPerGrid = TeamNum;
}
- INFO(DeviceId,
- "Launching kernel %s with %d blocks and %d threads in %s "
- "mode\n",
- (getOffloadEntry(DeviceId, TgtEntryPtr))
- ? getOffloadEntry(DeviceId, TgtEntryPtr)->name
- : "(null)",
- CudaBlocksPerGrid, CudaThreadsPerBlock,
- (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic");
+ if (getDebugLevel() || (getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL))
+ INFO(DeviceId,
+ "Launching kernel %s with %d blocks and %d threads in %s "
+ "mode\n",
+ (getOffloadEntry(DeviceId, TgtEntryPtr))
+ ? getOffloadEntry(DeviceId, TgtEntryPtr)->name
+ : "(null)",
+ CudaBlocksPerGrid, CudaThreadsPerBlock,
+ (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic");
CUstream Stream = getStream(DeviceId, AsyncInfo);
Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index 9d6f8bde1d0e..64a79f7a28c9 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -49,10 +49,11 @@ DeviceTy::DeviceTy(RTLInfoTy *RTL)
MemoryManager(nullptr) {}
DeviceTy::~DeviceTy() {
- if (DeviceID == -1 || getInfoLevel() < 1)
+ if (DeviceID == -1 || !(getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE))
return;
- dumpTargetPointerMappings(*this);
+ ident_t loc = {0, 0, 0, 0, ";libomptarget;libomptarget;0;0;;"};
+ dumpTargetPointerMappings(&loc, *this);
}
int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
@@ -217,14 +218,16 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
HT.incRefCount();
uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
- INFO(DeviceID,
- "Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
- ", "
- "Size=%" PRId64 ",%s RefCount=%s, Name=%s\n",
- (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(tp),
- Size, (UpdateRefCount ? " updated" : ""),
- HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str(),
- (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "(null)");
+ if (getDebugLevel() || getInfoLevel() & OMP_INFOTYPE_MAPPING_EXISTS)
+ INFO(DeviceID,
+ "Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
+ ", "
+ "Size=%" PRId64 ",%s RefCount=%s, Name=%s\n",
+ (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(tp),
+ Size, (UpdateRefCount ? " updated" : ""),
+ HT.isRefCountInf() ? "INF"
+ : std::to_string(HT.getRefCount()).c_str(),
+ (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
rc = (void *)tp;
} else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
// Explicit extension of mapped data - not allowed.
diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp
index b17be27275e7..b5af0b140585 100644
--- a/openmp/libomptarget/src/interface.cpp
+++ b/openmp/libomptarget/src/interface.cpp
@@ -57,22 +57,27 @@ static void HandleTargetOutcome(bool success, ident_t *loc = nullptr) {
break;
case tgt_mandatory:
if (!success) {
- if (getInfoLevel() > 1)
+ if (getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE)
for (const auto &Device : PM->Devices)
- dumpTargetPointerMappings(Device);
+ dumpTargetPointerMappings(loc, Device);
else
- FAILURE_MESSAGE("run with env LIBOMPTARGET_INFO>1 to dump host-target "
- "pointer maps\n");
+ FAILURE_MESSAGE("Run with LIBOMPTARGET_DEBUG=%d to dump host-target "
+ "pointer mappings.\n",
+ OMP_INFOTYPE_DUMP_TABLE);
SourceInfo info(loc);
if (info.isAvailible())
fprintf(stderr, "%s:%d:%d: ", info.getFilename(), info.getLine(),
info.getColumn());
else
- FAILURE_MESSAGE(
- "Build with debug information to provide more information");
+ FAILURE_MESSAGE("Source location information not present. Compile with "
+ "-g or -gline-tables-only.\n");
FATAL_MESSAGE0(
1, "failure of target construct while offloading is mandatory");
+ } else {
+ if (getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE)
+ for (const auto &Device : PM->Devices)
+ dumpTargetPointerMappings(loc, Device);
}
break;
}
@@ -147,12 +152,15 @@ EXTERN void __tgt_target_data_begin_mapper(ident_t *loc, int64_t device_id,
DeviceTy &Device = PM->Devices[device_id];
+ if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+ printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+ arg_names, "Entering OpenMP data region");
#ifdef OMPTARGET_DEBUG
for (int i = 0; i < arg_num; ++i) {
DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
", Type=0x%" PRIx64 ", Name=%s\n",
i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
- (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
+ (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#endif
@@ -227,12 +235,15 @@ EXTERN void __tgt_target_data_end_mapper(ident_t *loc, int64_t device_id,
return;
}
+ if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+ printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+ arg_names, "Exiting OpenMP data region");
#ifdef OMPTARGET_DEBUG
for (int i=0; i<arg_num; ++i) {
DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
", Type=0x%" PRIx64 ", Name=%s\n",
i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
- (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
+ (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#endif
@@ -294,6 +305,10 @@ EXTERN void __tgt_target_data_update_mapper(ident_t *loc, int64_t device_id,
return;
}
+ if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+ printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+ arg_names, "Updating OpenMP data");
+
DeviceTy &Device = PM->Devices[device_id];
int rc = targetDataUpdate(Device, arg_num, args_base, args, arg_sizes,
arg_types, arg_names, arg_mappers);
@@ -351,12 +366,15 @@ EXTERN int __tgt_target_mapper(ident_t *loc, int64_t device_id, void *host_ptr,
return OFFLOAD_FAIL;
}
+ if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+ printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+ arg_names, "Entering OpenMP kernel");
#ifdef OMPTARGET_DEBUG
for (int i=0; i<arg_num; ++i) {
DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
", Type=0x%" PRIx64 ", Name=%s\n",
i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
- (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
+ (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#endif
@@ -423,12 +441,15 @@ EXTERN int __tgt_target_teams_mapper(ident_t *loc, int64_t device_id,
return OFFLOAD_FAIL;
}
+ if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+ printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+ arg_names, "Entering OpenMP kernel");
#ifdef OMPTARGET_DEBUG
for (int i=0; i<arg_num; ++i) {
DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
", Type=0x%" PRIx64 ", Name=%s\n",
i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
- (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
+ (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#endif
@@ -436,7 +457,6 @@ EXTERN int __tgt_target_teams_mapper(ident_t *loc, int64_t device_id,
arg_types, arg_names, arg_mappers, team_num, thread_limit,
true /*team*/);
HandleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
-
return rc;
}
diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h
index a494e1ddb043..00826ae417dc 100644
--- a/openmp/libomptarget/src/private.h
+++ b/openmp/libomptarget/src/private.h
@@ -13,8 +13,9 @@
#ifndef _OMPTARGET_PRIVATE_H
#define _OMPTARGET_PRIVATE_H
-#include <omptarget.h>
#include <Debug.h>
+#include <SourceInfo.h>
+#include <omptarget.h>
#include <cstdint>
@@ -90,20 +91,60 @@ int __kmpc_get_target_offload(void) __attribute__((weak));
////////////////////////////////////////////////////////////////////////////////
/// dump a table of all the host-target pointer pairs on failure
-static inline void dumpTargetPointerMappings(const DeviceTy &Device) {
+static inline void dumpTargetPointerMappings(const ident_t *Loc,
+ const DeviceTy &Device) {
if (Device.HostDataToTargetMap.empty())
return;
- fprintf(stderr, "Device %d Host-Device Pointer Mappings:\n", Device.DeviceID);
- fprintf(stderr, "%-18s %-18s %s %s\n", "Host Ptr", "Target Ptr", "Size (B)",
- "Declaration");
+ SourceInfo Kernel(Loc);
+ INFO(Device.DeviceID,
+ "OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n",
+ Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn());
+ INFO(Device.DeviceID, "%-18s %-18s %s %s %s\n", "Host Ptr", "Target Ptr",
+ "Size (B)", "RefCount", "Declaration");
for (const auto &HostTargetMap : Device.HostDataToTargetMap) {
- SourceInfo info(HostTargetMap.HstPtrName);
- fprintf(stderr, DPxMOD " " DPxMOD " %-8lu %s at %s:%d:%d\n",
- DPxPTR(HostTargetMap.HstPtrBegin),
- DPxPTR(HostTargetMap.TgtPtrBegin),
- HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin, info.getName(),
- info.getFilename(), info.getLine(), info.getColumn());
+ SourceInfo Info(HostTargetMap.HstPtrName);
+ INFO(Device.DeviceID, DPxMOD " " DPxMOD " %-8lu %-8ld %s at %s:%d:%d\n",
+ DPxPTR(HostTargetMap.HstPtrBegin), DPxPTR(HostTargetMap.TgtPtrBegin),
+ (long unsigned)(HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin),
+ HostTargetMap.getRefCount(), Info.getName(), Info.getFilename(),
+ Info.getLine(), Info.getColumn());
+ }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Print out the names and properties of the arguments to each kernel
+static inline void
+printKernelArguments(const ident_t *Loc, const int64_t DeviceId,
+ const int32_t ArgNum, const int64_t *ArgSizes,
+ const int64_t *ArgTypes, const map_var_info_t *ArgNames,
+ const char *RegionType) {
+ SourceInfo info(Loc);
+ INFO(DeviceId, "%s at %s:%d:%d with %d arguments:\n", RegionType,
+ info.getFilename(), info.getLine(), info.getColumn(), ArgNum);
+
+ for (int32_t i = 0; i < ArgNum; ++i) {
+ const map_var_info_t varName = (ArgNames) ? ArgNames[i] : nullptr;
+ const char *type = nullptr;
+ const char *implicit =
+ (ArgTypes[i] & OMP_TGT_MAPTYPE_IMPLICIT) ? "(implicit)" : "";
+ if (ArgTypes[i] & OMP_TGT_MAPTYPE_TO && ArgTypes[i] & OMP_TGT_MAPTYPE_FROM)
+ type = "tofrom";
+ else if (ArgTypes[i] & OMP_TGT_MAPTYPE_TO)
+ type = "to";
+ else if (ArgTypes[i] & OMP_TGT_MAPTYPE_FROM)
+ type = "from";
+ else if (ArgTypes[i] & OMP_TGT_MAPTYPE_PRIVATE)
+ type = "private";
+ else if (ArgTypes[i] & OMP_TGT_MAPTYPE_LITERAL)
+ type = "firstprivate";
+ else if (ArgTypes[i] & OMP_TGT_MAPTYPE_TARGET_PARAM && ArgSizes[i] != 0)
+ type = "alloc";
+ else
+ type = "use_address";
+
+ INFO(DeviceId, "%s(%s)[%ld] %s\n", type,
+ getNameFromMapping(varName).c_str(), ArgSizes[i], implicit);
}
}
diff --git a/openmp/libomptarget/test/offloading/info.c b/openmp/libomptarget/test/offloading/info.c
index e0d3f1a0e94c..e04f9ccaaf42 100644
--- a/openmp/libomptarget/test/offloading/info.c
+++ b/openmp/libomptarget/test/offloading/info.c
@@ -1,15 +1,38 @@
-// RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_INFO=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -gline-tables-only && env LIBOMPTARGET_INFO=23 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
#include <stdio.h>
#include <omp.h>
+#define N 64
+
int main() {
- int ptr = 1;
+ int A[N];
+ int B[N];
+ int C[N];
+ int val = 1;
-// INFO: CUDA device {{[0-9]+}} info: Device supports up to {{[0-9]+}} CUDA blocks and {{[0-9]+}} threads with a warp size of {{[0-9]+}}
-// INFO: CUDA device {{[0-9]+}} info: Launching kernel {{.*}} with {{[0-9]+}} blocks and {{[0-9]+}} threads in Generic mode
-#pragma omp target map(tofrom:ptr)
- {ptr = 1;}
+// INFO: CUDA device 0 info: Device supports up to {{.*}} CUDA blocks and {{.*}} threads with a warp size of {{.*}}
+// INFO: Libomptarget device 0 info: Entering OpenMP data region at info.c:33:1 with 3 arguments:
+// INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
+// INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
+// INFO: Libomptarget device 0 info: to(C[0:64])[256]
+// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:33:1:
+// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 C[0:64] at info.c:11:7
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 B[0:64] at info.c:10:7
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 A[0:64] at info.c:9:7
+// INFO: Libomptarget device 0 info: Entering OpenMP kernel at info.c:34:1 with 1 arguments:
+// INFO: Libomptarget device 0 info: firstprivate(val)[4]
+// INFO: CUDA device 0 info: Launching kernel {{.*}} with {{.*}} and {{.*}} threads in {{.*}} mode
+// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:34:1:
+// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256 1 C[0:64] at info.c:11:7
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256 1 B[0:64] at info.c:10:7
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256 1 A[0:64] at info.c:9:7
+// INFO: Libomptarget device 0 info: Exiting OpenMP data region at info.c:33:1
+#pragma omp target data map(alloc:A[0:N]) map(tofrom:B[0:N]) map(to:C[0:N])
+#pragma omp target firstprivate(val)
+ { val = 1; }
return 0;
}
More information about the llvm-branch-commits
mailing list