[Openmp-commits] [openmp] [Libomptarget][NFC] Format in-line comments consistently (PR #77530)
Joseph Huber via Openmp-commits
openmp-commits at lists.llvm.org
Wed Jan 10 08:10:11 PST 2024
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/77530
>From 452744131e8d353fa556cd2aebf11f425b564aef Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Tue, 9 Jan 2024 15:54:48 -0600
Subject: [PATCH] [Libomptarget][NFC] Format in-line comments consistently
Summary:
The LLVM style uses /*Foo=*/ when indicating the name of a constant. See
https://llvm.org/docs/CodingStandards.html#comment-formatting. This is
useful for consistency, as well as because `clang-format` understands
this syntax and formats it more cleanly. Do a bulk update of this
syntax.
---
openmp/libomptarget/DeviceRTL/include/State.h | 26 +++++-----
openmp/libomptarget/DeviceRTL/src/Kernel.cpp | 4 +-
.../DeviceRTL/src/Parallelism.cpp | 26 +++++-----
.../libomptarget/DeviceRTL/src/Reduction.cpp | 2 +-
openmp/libomptarget/include/Shared/Profile.h | 2 +-
.../plugins-nextgen/amdgpu/src/rtl.cpp | 14 +++---
.../common/include/GlobalHandler.h | 8 ++--
.../plugins-nextgen/common/src/JIT.cpp | 8 ++--
.../common/src/PluginInterface.cpp | 48 ++++++++-----------
.../plugins-nextgen/cuda/src/rtl.cpp | 8 ++--
.../generic-elf-64bit/src/rtl.cpp | 4 +-
openmp/libomptarget/src/OpenMP/Mapping.cpp | 18 +++----
.../libomptarget/src/OpenMP/OMPT/Callback.cpp | 36 +++++++-------
openmp/libomptarget/src/device.cpp | 8 ++--
openmp/libomptarget/src/interface.cpp | 6 +--
openmp/libomptarget/src/omptarget.cpp | 6 +--
.../kernelreplay/llvm-omp-kernel-replay.cpp | 24 +++++-----
17 files changed, 120 insertions(+), 128 deletions(-)
diff --git a/openmp/libomptarget/DeviceRTL/include/State.h b/openmp/libomptarget/DeviceRTL/include/State.h
index c93de4191f83fa..1a3490394458ff 100644
--- a/openmp/libomptarget/DeviceRTL/include/State.h
+++ b/openmp/libomptarget/DeviceRTL/include/State.h
@@ -240,29 +240,29 @@ lookupPtr(ValueKind Kind, bool IsReadonly, bool ForceTeamState) {
/// update ICV values we can declare in global scope.
template <typename Ty, ValueKind Kind> struct Value {
[[gnu::flatten, gnu::always_inline]] operator Ty() {
- return lookup(/* IsReadonly */ true, /* IdentTy */ nullptr,
- /* ForceTeamState */ false);
+ return lookup(/*IsReadonly=*/true, /*IdentTy=*/nullptr,
+ /*ForceTeamState=*/false);
}
[[gnu::flatten, gnu::always_inline]] Value &operator=(const Ty &Other) {
- set(Other, /* IdentTy */ nullptr);
+ set(Other, /*IdentTy=*/nullptr);
return *this;
}
[[gnu::flatten, gnu::always_inline]] Value &operator++() {
- inc(1, /* IdentTy */ nullptr);
+ inc(1, /*IdentTy=*/nullptr);
return *this;
}
[[gnu::flatten, gnu::always_inline]] Value &operator--() {
- inc(-1, /* IdentTy */ nullptr);
+ inc(-1, /*IdentTy=*/nullptr);
return *this;
}
[[gnu::flatten, gnu::always_inline]] void
assert_eq(const Ty &V, IdentTy *Ident = nullptr,
bool ForceTeamState = false) {
- ASSERT(lookup(/* IsReadonly */ true, Ident, ForceTeamState) == V, nullptr);
+ ASSERT(lookup(/*IsReadonly=*/true, Ident, ForceTeamState) == V, nullptr);
}
private:
@@ -273,12 +273,12 @@ template <typename Ty, ValueKind Kind> struct Value {
}
[[gnu::flatten, gnu::always_inline]] Ty &inc(int UpdateVal, IdentTy *Ident) {
- return (lookup(/* IsReadonly */ false, Ident, /* ForceTeamState */ false) +=
+ return (lookup(/*IsReadonly=*/false, Ident, /*ForceTeamState=*/false) +=
UpdateVal);
}
[[gnu::flatten, gnu::always_inline]] Ty &set(Ty UpdateVal, IdentTy *Ident) {
- return (lookup(/* IsReadonly */ false, Ident, /* ForceTeamState */ false) =
+ return (lookup(/*IsReadonly=*/false, Ident, /*ForceTeamState=*/false) =
UpdateVal);
}
@@ -290,8 +290,8 @@ template <typename Ty, ValueKind Kind> struct Value {
/// we can declare in global scope.
template <typename Ty, ValueKind Kind> struct PtrValue {
[[gnu::flatten, gnu::always_inline]] operator Ty() {
- return lookup(/* IsReadonly */ true, /* IdentTy */ nullptr,
- /* ForceTeamState */ false);
+ return lookup(/*IsReadonly=*/true, /*IdentTy=*/nullptr,
+ /*ForceTeamState=*/false);
}
[[gnu::flatten, gnu::always_inline]] PtrValue &operator=(const Ty Other) {
@@ -305,8 +305,8 @@ template <typename Ty, ValueKind Kind> struct PtrValue {
}
Ty &set(Ty UpdateVal) {
- return (lookup(/* IsReadonly */ false, /* IdentTy */ nullptr,
- /* ForceTeamState */ false) = UpdateVal);
+ return (lookup(/*IsReadonly=*/false, /*IdentTy=*/nullptr,
+ /*ForceTeamState=*/false) = UpdateVal);
}
template <typename VTy, typename Ty2> friend struct ValueRAII;
@@ -315,7 +315,7 @@ template <typename Ty, ValueKind Kind> struct PtrValue {
template <typename VTy, typename Ty> struct ValueRAII {
ValueRAII(VTy &V, Ty NewValue, Ty OldValue, bool Active, IdentTy *Ident,
bool ForceTeamState = false)
- : Ptr(Active ? &V.lookup(/* IsReadonly */ false, Ident, ForceTeamState)
+ : Ptr(Active ? &V.lookup(/*IsReadonly=*/false, Ident, ForceTeamState)
: (Ty *)utils::UndefPtr),
Val(OldValue), Active(Active) {
if (!Active)
diff --git a/openmp/libomptarget/DeviceRTL/src/Kernel.cpp b/openmp/libomptarget/DeviceRTL/src/Kernel.cpp
index 06b12fec21677c..95d4c728016d24 100644
--- a/openmp/libomptarget/DeviceRTL/src/Kernel.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Kernel.cpp
@@ -78,11 +78,11 @@ int32_t __kmpc_target_init(KernelEnvironmentTy &KernelEnvironment,
llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD;
bool UseGenericStateMachine = Configuration.UseGenericStateMachine;
if (IsSPMD) {
- inititializeRuntime(/* IsSPMD */ true, KernelEnvironment,
+ inititializeRuntime(/*IsSPMD=*/true, KernelEnvironment,
KernelLaunchEnvironment);
synchronize::threadsAligned(atomic::relaxed);
} else {
- inititializeRuntime(/* IsSPMD */ false, KernelEnvironment,
+ inititializeRuntime(/*IsSPMD=*/false, KernelEnvironment,
KernelLaunchEnvironment);
// No need to wait since only the main threads will execute user
// code and workers will run into a barrier right away.
diff --git a/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp b/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp
index 2c0701bd5358fd..7005477bf4c79d 100644
--- a/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp
@@ -121,20 +121,20 @@ __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr,
// created.
state::ValueRAII ParallelTeamSizeRAII(state::ParallelTeamSize, PTeamSize,
1u, TId == 0, ident,
- /* ForceTeamState */ true);
+ /*ForceTeamState=*/true);
state::ValueRAII ActiveLevelRAII(icv::ActiveLevel, 1u, 0u, TId == 0,
- ident, /* ForceTeamState */ true);
+ ident, /*ForceTeamState=*/true);
state::ValueRAII LevelRAII(icv::Level, 1u, 0u, TId == 0, ident,
- /* ForceTeamState */ true);
+ /*ForceTeamState=*/true);
// Synchronize all threads after the main thread (TId == 0) set up the
// team state properly.
synchronize::threadsAligned(atomic::acq_rel);
state::ParallelTeamSize.assert_eq(PTeamSize, ident,
- /* ForceTeamState */ true);
- icv::ActiveLevel.assert_eq(1u, ident, /* ForceTeamState */ true);
- icv::Level.assert_eq(1u, ident, /* ForceTeamState */ true);
+ /*ForceTeamState=*/true);
+ icv::ActiveLevel.assert_eq(1u, ident, /*ForceTeamState=*/true);
+ icv::Level.assert_eq(1u, ident, /*ForceTeamState=*/true);
// Ensure we synchronize before we run user code to avoid invalidating the
// assumptions above.
@@ -152,9 +152,9 @@ __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr,
// __kmpc_target_deinit may not hold.
synchronize::threadsAligned(atomic::acq_rel);
- state::ParallelTeamSize.assert_eq(1u, ident, /* ForceTeamState */ true);
- icv::ActiveLevel.assert_eq(0u, ident, /* ForceTeamState */ true);
- icv::Level.assert_eq(0u, ident, /* ForceTeamState */ true);
+ state::ParallelTeamSize.assert_eq(1u, ident, /*ForceTeamState=*/true);
+ icv::ActiveLevel.assert_eq(0u, ident, /*ForceTeamState=*/true);
+ icv::Level.assert_eq(0u, ident, /*ForceTeamState=*/true);
// Ensure we synchronize to create an aligned region around the assumptions.
synchronize::threadsAligned(atomic::relaxed);
@@ -242,14 +242,14 @@ __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr,
// created.
state::ValueRAII ParallelTeamSizeRAII(state::ParallelTeamSize, PTeamSize,
1u, true, ident,
- /* ForceTeamState */ true);
+ /*ForceTeamState=*/true);
state::ValueRAII ParallelRegionFnRAII(state::ParallelRegionFn, wrapper_fn,
(void *)nullptr, true, ident,
- /* ForceTeamState */ true);
+ /*ForceTeamState=*/true);
state::ValueRAII ActiveLevelRAII(icv::ActiveLevel, 1u, 0u, true, ident,
- /* ForceTeamState */ true);
+ /*ForceTeamState=*/true);
state::ValueRAII LevelRAII(icv::Level, 1u, 0u, true, ident,
- /* ForceTeamState */ true);
+ /*ForceTeamState=*/true);
// Master signals work to activate workers.
synchronize::threads(atomic::seq_cst);
diff --git a/openmp/libomptarget/DeviceRTL/src/Reduction.cpp b/openmp/libomptarget/DeviceRTL/src/Reduction.cpp
index 0a9db3edabccbf..744d1a3a231c83 100644
--- a/openmp/libomptarget/DeviceRTL/src/Reduction.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Reduction.cpp
@@ -69,7 +69,7 @@ static int32_t nvptx_parallel_reduce_nowait(void *reduce_data,
ShuffleReductFnTy shflFct,
InterWarpCopyFnTy cpyFct) {
uint32_t BlockThreadId = mapping::getThreadIdInBlock();
- if (mapping::isMainThreadInGenericMode(/* IsSPMD */ false))
+ if (mapping::isMainThreadInGenericMode(/*IsSPMD=*/false))
BlockThreadId = 0;
uint32_t NumThreads = omp_get_num_threads();
if (NumThreads == 1)
diff --git a/openmp/libomptarget/include/Shared/Profile.h b/openmp/libomptarget/include/Shared/Profile.h
index 7e580988a39baf..39817bab36cbb2 100644
--- a/openmp/libomptarget/include/Shared/Profile.h
+++ b/openmp/libomptarget/include/Shared/Profile.h
@@ -32,7 +32,7 @@ class Profiler {
Int32Envar ProfileGranularity =
Int32Envar("LIBOMPTARGET_PROFILE_GRANULARITY", 500);
- llvm::timeTraceProfilerInitialize(ProfileGranularity /* us */,
+ llvm::timeTraceProfilerInitialize(ProfileGranularity /*us=*/,
"libomptarget");
}
diff --git a/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp
index b67642e9e1bcb3..8424d0f5df0822 100644
--- a/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -2211,10 +2211,9 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
hsa_amd_pointer_info_t Info;
Info.size = sizeof(hsa_amd_pointer_info_t);
- hsa_status_t Status =
- hsa_amd_pointer_info(HstPtr, &Info, /* Allocator */ nullptr,
- /* Number of accessible agents (out) */ nullptr,
- /* Accessible agents */ nullptr);
+ hsa_status_t Status = hsa_amd_pointer_info(
+ HstPtr, &Info, /*Allocator=*/nullptr, /*num_agents_accessible=*/nullptr,
+ /*accessible=*/nullptr);
if (auto Err = Plugin::check(Status, "Error in hsa_amd_pointer_info: %s"))
return std::move(Err);
@@ -2789,7 +2788,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
AMDHostDeviceTy &HostDevice;
/// The current size of the global device memory pool (managed by us).
- uint64_t DeviceMemoryPoolSize = 1L << 29L /* 512MB */;
+ uint64_t DeviceMemoryPoolSize = 1L << 29L /*512MB=*/;
/// The current size of the stack that will be used in cases where it could
/// not be statically determined.
@@ -3031,9 +3030,8 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
/// Check whether the image is compatible with an AMDGPU device.
Expected<bool> isELFCompatible(StringRef Image) const override {
// Get the associated architecture and flags from the ELF.
- auto ElfOrErr =
- ELF64LEObjectFile::create(MemoryBufferRef(Image, /*Identifier=*/""),
- /*InitContent=*/false);
+ auto ElfOrErr = ELF64LEObjectFile::create(
+ MemoryBufferRef(Image, /*Identifier=*/""), /*InitContent=*/false);
if (!ElfOrErr)
return ElfOrErr.takeError();
std::optional<StringRef> Processor = ElfOrErr->tryGetCPUName();
diff --git a/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h b/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h
index d9fe938790ca76..8707e7b4c504e3 100644
--- a/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h
+++ b/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h
@@ -138,7 +138,7 @@ class GenericGlobalHandlerTy {
const GlobalTy &HostGlobal,
const GlobalTy &DeviceGlobal) {
return moveGlobalBetweenDeviceAndHost(Device, HostGlobal, DeviceGlobal,
- /* D2H */ true);
+ /*D2H=*/true);
}
/// Copy the memory associated with a global from the device to its
@@ -147,7 +147,7 @@ class GenericGlobalHandlerTy {
Error readGlobalFromDevice(GenericDeviceTy &Device, DeviceImageTy &Image,
const GlobalTy &HostGlobal) {
return moveGlobalBetweenDeviceAndHost(Device, Image, HostGlobal,
- /* D2H */ true);
+ /*D2H=*/true);
}
/// Copy the memory associated with a global from the host to its counterpart
@@ -156,7 +156,7 @@ class GenericGlobalHandlerTy {
Error writeGlobalToDevice(GenericDeviceTy &Device, const GlobalTy &HostGlobal,
const GlobalTy &DeviceGlobal) {
return moveGlobalBetweenDeviceAndHost(Device, HostGlobal, DeviceGlobal,
- /* D2H */ false);
+ /*D2H=*/false);
}
/// Copy the memory associated with a global from the host to its counterpart
@@ -165,7 +165,7 @@ class GenericGlobalHandlerTy {
Error writeGlobalToDevice(GenericDeviceTy &Device, DeviceImageTy &Image,
const GlobalTy &HostGlobal) {
return moveGlobalBetweenDeviceAndHost(Device, Image, HostGlobal,
- /* D2H */ false);
+ /*D2H=*/false);
}
};
diff --git a/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp b/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp
index 7275be4edfca5b..9eb610cab4de66 100644
--- a/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp
+++ b/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp
@@ -93,7 +93,7 @@ createModuleFromImage(const __tgt_device_image &Image, LLVMContext &Context) {
StringRef Data((const char *)Image.ImageStart,
target::getPtrDiff(Image.ImageEnd, Image.ImageStart));
std::unique_ptr<MemoryBuffer> MB = MemoryBuffer::getMemBuffer(
- Data, /* BufferName */ "", /* RequiresNullTerminator */ false);
+ Data, /*BufferName=*/"", /*RequiresNullTerminator=*/false);
return createModuleFromMemoryBuffer(MB, Context);
}
@@ -186,7 +186,7 @@ void JITEngine::codegen(TargetMachine *TM, TargetLibraryInfoImpl *TLII,
TM->addPassesToEmitFile(PM, OS, nullptr,
TT.isNVPTX() ? CodeGenFileType::AssemblyFile
: CodeGenFileType::ObjectFile,
- /* DisableVerify */ false, MMIWP);
+ /*DisableVerify=*/false, MMIWP);
PM.run(M);
}
@@ -196,8 +196,8 @@ JITEngine::backend(Module &M, const std::string &ComputeUnitKind,
unsigned OptLevel) {
auto RemarksFileOrErr = setupLLVMOptimizationRemarks(
- M.getContext(), /* RemarksFilename */ "", /* RemarksPasses */ "",
- /* RemarksFormat */ "", /* RemarksWithHotness */ false);
+ M.getContext(), /*RemarksFilename=*/"", /*RemarksPasses=*/"",
+ /*RemarksFormat=*/"", /*RemarksWithHotness=*/false);
if (Error E = RemarksFileOrErr.takeError())
return std::move(E);
if (*RemarksFileOrErr)
diff --git a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
index 9490e58fc669cd..0db7910ec10567 100644
--- a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
@@ -64,7 +64,7 @@ struct RecordReplayTy {
void *suggestAddress(uint64_t MaxMemoryAllocation) {
// Get a valid pointer address for this system
void *Addr =
- Device->allocate(1024, /* HstPtr */ nullptr, TARGET_ALLOC_DEFAULT);
+ Device->allocate(1024, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT);
Device->free(Addr);
// Align Address to MaxMemoryAllocation
Addr = (void *)alignPtr((Addr), MaxMemoryAllocation);
@@ -104,8 +104,8 @@ struct RecordReplayTy {
constexpr size_t STEP = 1024 * 1024 * 1024ULL;
MemoryStart = nullptr;
for (TotalSize = MAX_MEMORY_ALLOCATION; TotalSize > 0; TotalSize -= STEP) {
- MemoryStart = Device->allocate(TotalSize, /* HstPtr */ nullptr,
- TARGET_ALLOC_DEFAULT);
+ MemoryStart =
+ Device->allocate(TotalSize, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT);
if (MemoryStart)
break;
}
@@ -214,8 +214,8 @@ struct RecordReplayTy {
for (auto &OffloadEntry : Image.getOffloadEntryTable()) {
if (!OffloadEntry.size)
continue;
- Size += std::strlen(OffloadEntry.name) + /* '\0' */ 1 +
- /* OffloadEntry.size value */ sizeof(uint32_t) +
+ // Get the total size of the string and entry including the null byte.
+ Size += std::strlen(OffloadEntry.name) + 1 + sizeof(uint32_t) +
OffloadEntry.size;
}
@@ -735,13 +735,12 @@ Error GenericDeviceTy::init(GenericPluginTy &Plugin) {
if (ompt::Initialized) {
bool ExpectedStatus = false;
if (OmptInitialized.compare_exchange_strong(ExpectedStatus, true))
- performOmptCallback(device_initialize,
- /* device_num */ DeviceId +
- Plugin.getDeviceIdStartIndex(),
- /* type */ getComputeUnitKind().c_str(),
- /* device */ reinterpret_cast<ompt_device_t *>(this),
- /* lookup */ ompt::lookupCallbackByName,
- /* documentation */ nullptr);
+ performOmptCallback(device_initialize, /*device_num=*/DeviceId +
+ Plugin.getDeviceIdStartIndex(),
+ /*type=*/getComputeUnitKind().c_str(),
+ /*device=*/reinterpret_cast<ompt_device_t *>(this),
+ /*lookup=*/ompt::lookupCallbackByName,
+ /*documentation=*/nullptr);
}
#endif
@@ -835,7 +834,7 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
bool ExpectedStatus = true;
if (OmptInitialized.compare_exchange_strong(ExpectedStatus, false))
performOmptCallback(device_finalize,
- /* device_num */ DeviceId +
+ /*device_num=*/DeviceId +
Plugin.getDeviceIdStartIndex());
}
#endif
@@ -897,16 +896,11 @@ GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,
if (ompt::Initialized) {
size_t Bytes =
getPtrDiff(InputTgtImage->ImageEnd, InputTgtImage->ImageStart);
- performOmptCallback(device_load,
- /* device_num */ DeviceId +
- Plugin.getDeviceIdStartIndex(),
- /* FileName */ nullptr,
- /* File Offset */ 0,
- /* VmaInFile */ nullptr,
- /* ImgSize */ Bytes,
- /* HostAddr */ InputTgtImage->ImageStart,
- /* DeviceAddr */ nullptr,
- /* FIXME: ModuleId */ 0);
+ performOmptCallback(
+ device_load, /*device_num=*/DeviceId + Plugin.getDeviceIdStartIndex(),
+ /*FileName=*/nullptr, /*FileOffset=*/0, /*VmaInFile=*/nullptr,
+ /*ImgSize=*/Bytes, /*HostAddr=*/InputTgtImage->ImageStart,
+ /*DeviceAddr=*/nullptr, /* FIXME: ModuleId */ 0);
}
#endif
@@ -1293,7 +1287,7 @@ Error PinnedAllocationMapTy::lockMappedHostBuffer(void *HstPtr, size_t Size) {
// If pinned, just insert the entry representing the whole pinned buffer.
if (*IsPinnedOrErr)
return insertEntry(BaseHstPtr, BaseDevAccessiblePtr, BaseSize,
- /* Externally locked */ true);
+ /*Externallylocked=*/true);
// Not externally pinned. Do nothing if locking of mapped buffers is disabled.
if (!LockMappedBuffers)
@@ -1863,7 +1857,7 @@ int32_t __tgt_rtl_data_notify_unmapped(int32_t DeviceId, void *HstPtr) {
int32_t __tgt_rtl_data_submit(int32_t DeviceId, void *TgtPtr, void *HstPtr,
int64_t Size) {
return __tgt_rtl_data_submit_async(DeviceId, TgtPtr, HstPtr, Size,
- /* AsyncInfoPtr */ nullptr);
+ /*AsyncInfoPtr=*/nullptr);
}
int32_t __tgt_rtl_data_submit_async(int32_t DeviceId, void *TgtPtr,
@@ -1885,7 +1879,7 @@ int32_t __tgt_rtl_data_submit_async(int32_t DeviceId, void *TgtPtr,
int32_t __tgt_rtl_data_retrieve(int32_t DeviceId, void *HstPtr, void *TgtPtr,
int64_t Size) {
return __tgt_rtl_data_retrieve_async(DeviceId, HstPtr, TgtPtr, Size,
- /* AsyncInfoPtr */ nullptr);
+ /*AsyncInfoPtr=*/nullptr);
}
int32_t __tgt_rtl_data_retrieve_async(int32_t DeviceId, void *HstPtr,
@@ -1909,7 +1903,7 @@ int32_t __tgt_rtl_data_exchange(int32_t SrcDeviceId, void *SrcPtr,
int64_t Size) {
return __tgt_rtl_data_exchange_async(SrcDeviceId, SrcPtr, DstDeviceId, DstPtr,
Size,
- /* AsyncInfoPtr */ nullptr);
+ /*AsyncInfoPtr=*/nullptr);
}
int32_t __tgt_rtl_data_exchange_async(int32_t SrcDeviceId, void *SrcPtr,
diff --git a/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp
index 0005bff7a8035d..fb51f96c07b74d 100644
--- a/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp
+++ b/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp
@@ -1216,10 +1216,10 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
std::max(KernelArgs.DynCGroupMem, GenericDevice.getDynamicMemorySize());
CUresult Res =
- cuLaunchKernel(Func, NumBlocks, /* gridDimY */ 1,
- /* gridDimZ */ 1, NumThreads,
- /* blockDimY */ 1, /* blockDimZ */ 1, MaxDynCGroupMem,
- Stream, (void **)Args, nullptr);
+ cuLaunchKernel(Func, NumBlocks, /*gridDimY=*/1,
+ /*gridDimZ=*/1, NumThreads,
+ /*blockDimY=*/1, /*blockDimZ=*/1, MaxDynCGroupMem, Stream,
+ (void **)Args, nullptr);
return Plugin::check(Res, "Error in cuLaunchKernel for '%s': %s", getName());
}
diff --git a/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp
index 7f66b6827bce00..6466afc543b568 100644
--- a/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp
+++ b/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp
@@ -73,8 +73,8 @@ struct GenELF64KernelTy : public GenericKernelTy {
Func = (void (*)())Global.getPtr();
KernelEnvironment.Configuration.ExecMode = OMP_TGT_EXEC_MODE_GENERIC;
- KernelEnvironment.Configuration.MayUseNestedParallelism = /* Unknown */ 2;
- KernelEnvironment.Configuration.UseGenericStateMachine = /* Unknown */ 2;
+ KernelEnvironment.Configuration.MayUseNestedParallelism = /*Unknown=*/2;
+ KernelEnvironment.Configuration.UseGenericStateMachine = /*Unknown=*/2;
// Set the maximum number of threads to a single.
MaxNumThreads = 1;
diff --git a/openmp/libomptarget/src/OpenMP/Mapping.cpp b/openmp/libomptarget/src/OpenMP/Mapping.cpp
index a5c24810e0af95..833856f2abf290 100644
--- a/openmp/libomptarget/src/OpenMP/Mapping.cpp
+++ b/openmp/libomptarget/src/OpenMP/Mapping.cpp
@@ -303,9 +303,9 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// Notify the plugin about the new mapping.
if (Device.notifyDataMapped(HstPtrBegin, Size))
- return {{false /* IsNewEntry */, false /* IsHostPointer */},
- nullptr /* Entry */,
- nullptr /* TargetPointer */};
+ return {{false /*IsNewEntry=*/, false /*IsHostPointer=*/},
+ nullptr /*Entry=*/,
+ nullptr /*TargetPointer=*/};
} else {
// This entry is not present and we did not create a new entry for it.
LR.TPR.Flags.IsPresent = false;
@@ -333,9 +333,9 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
LR.TPR.TargetPointer = nullptr;
} else if (LR.TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
OFFLOAD_SUCCESS)
- return {{false /* IsNewEntry */, false /* IsHostPointer */},
- nullptr /* Entry */,
- nullptr /* TargetPointer */};
+ return {{false /*IsNewEntry=*/, false /*IsHostPointer=*/},
+ nullptr /*Entry=*/,
+ nullptr /*TargetPointer=*/};
} else {
// If not a host pointer and no present modifier, we need to wait for the
// event if it exists.
@@ -349,9 +349,9 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// If it fails to wait for the event, we need to return nullptr in
// case of any data race.
REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
- return {{false /* IsNewEntry */, false /* IsHostPointer */},
- nullptr /* Entry */,
- nullptr /* TargetPointer */};
+ return {{false /*IsNewEntry=*/, false /*IsHostPointer=*/},
+ nullptr /*Entry=*/,
+ nullptr /*TargetPointer=*/};
}
}
}
diff --git a/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp b/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp
index da955e101956f4..82934f10486c5a 100644
--- a/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp
+++ b/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp
@@ -88,16 +88,16 @@ void Interface::beginTargetDataAlloc(int64_t DeviceId, void *HstPtrBegin,
ompt_callback_target_data_op_emi_fn(
ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_alloc, HstPtrBegin,
- /* SrcDeviceNum */ omp_get_initial_device(), *TgtPtrBegin,
- /* TgtDeviceNum */ DeviceId, Size, Code);
+ /*SrcDeviceNum=*/omp_get_initial_device(), *TgtPtrBegin,
+ /*TgtDeviceNum=*/DeviceId, Size, Code);
} else if (ompt_callback_target_data_op_fn) {
// HostOpId is set by the runtime
HostOpId = createOpId();
// Invoke the tool supplied data op callback
ompt_callback_target_data_op_fn(
TargetData.value, HostOpId, ompt_target_data_alloc, HstPtrBegin,
- /* SrcDeviceNum */ omp_get_initial_device(), *TgtPtrBegin,
- /* TgtDeviceNum */ DeviceId, Size, Code);
+ /*SrcDeviceNum=*/omp_get_initial_device(), *TgtPtrBegin,
+ /*TgtDeviceNum=*/DeviceId, Size, Code);
}
}
@@ -111,8 +111,8 @@ void Interface::endTargetDataAlloc(int64_t DeviceId, void *HstPtrBegin,
ompt_callback_target_data_op_emi_fn(
ompt_scope_end, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_alloc, HstPtrBegin,
- /* SrcDeviceNum */ omp_get_initial_device(), *TgtPtrBegin,
- /* TgtDeviceNum */ DeviceId, Size, Code);
+ /*SrcDeviceNum=*/omp_get_initial_device(), *TgtPtrBegin,
+ /*TgtDeviceNum=*/DeviceId, Size, Code);
}
endTargetDataOperation();
}
@@ -127,15 +127,15 @@ void Interface::beginTargetDataSubmit(int64_t DeviceId, void *TgtPtrBegin,
ompt_callback_target_data_op_emi_fn(
ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_transfer_to_device, HstPtrBegin,
- /* SrcDeviceNum */ omp_get_initial_device(), TgtPtrBegin, DeviceId,
- Size, Code);
+ /*SrcDeviceNum=*/omp_get_initial_device(), TgtPtrBegin, DeviceId, Size,
+ Code);
} else if (ompt_callback_target_data_op_fn) {
// HostOpId is set by the runtime
HostOpId = createOpId();
// Invoke the tool supplied data op callback
ompt_callback_target_data_op_fn(
TargetData.value, HostOpId, ompt_target_data_transfer_to_device,
- HstPtrBegin, /* SrcDeviceNum */ omp_get_initial_device(), TgtPtrBegin,
+ HstPtrBegin, /*SrcDeviceNum=*/omp_get_initial_device(), TgtPtrBegin,
DeviceId, Size, Code);
}
}
@@ -150,8 +150,8 @@ void Interface::endTargetDataSubmit(int64_t DeviceId, void *TgtPtrBegin,
ompt_callback_target_data_op_emi_fn(
ompt_scope_end, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_transfer_to_device, HstPtrBegin,
- /* SrcDeviceNum */ omp_get_initial_device(), TgtPtrBegin, DeviceId,
- Size, Code);
+ /*SrcDeviceNum=*/omp_get_initial_device(), TgtPtrBegin, DeviceId, Size,
+ Code);
}
endTargetDataOperation();
}
@@ -165,15 +165,15 @@ void Interface::beginTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin,
ompt_callback_target_data_op_emi_fn(
ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_delete, TgtPtrBegin, DeviceId,
- /* TgtPtrBegin */ nullptr, /* TgtDeviceNum */ -1, /* Bytes */ 0, Code);
+ /*TgtPtrBegin=*/nullptr, /*TgtDeviceNum=*/-1, /*Bytes=*/0, Code);
} else if (ompt_callback_target_data_op_fn) {
// HostOpId is set by the runtime
HostOpId = createOpId();
// Invoke the tool supplied data op callback
ompt_callback_target_data_op_fn(TargetData.value, HostOpId,
ompt_target_data_delete, TgtPtrBegin,
- DeviceId, /* TgtPtrBegin */ nullptr,
- /* TgtDeviceNum */ -1, /* Bytes */ 0, Code);
+ DeviceId, /*TgtPtrBegin=*/nullptr,
+ /*TgtDeviceNum=*/-1, /*Bytes=*/0, Code);
}
}
@@ -186,7 +186,7 @@ void Interface::endTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin,
ompt_callback_target_data_op_emi_fn(
ompt_scope_end, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_delete, TgtPtrBegin, DeviceId,
- /* TgtPtrBegin */ nullptr, /* TgtDeviceNum */ -1, /* Bytes */ 0, Code);
+ /*TgtPtrBegin=*/nullptr, /*TgtDeviceNum=*/-1, /*Bytes=*/0, Code);
}
endTargetDataOperation();
}
@@ -202,7 +202,7 @@ void Interface::beginTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin,
ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_transfer_from_device, TgtPtrBegin, DeviceId,
HstPtrBegin,
- /* TgtDeviceNum */ omp_get_initial_device(), Size, Code);
+ /*TgtDeviceNum=*/omp_get_initial_device(), Size, Code);
} else if (ompt_callback_target_data_op_fn) {
// HostOpId is set by the runtime
HostOpId = createOpId();
@@ -210,7 +210,7 @@ void Interface::beginTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin,
ompt_callback_target_data_op_fn(
TargetData.value, HostOpId, ompt_target_data_transfer_from_device,
TgtPtrBegin, DeviceId, HstPtrBegin,
- /* TgtDeviceNum */ omp_get_initial_device(), Size, Code);
+ /*TgtDeviceNum=*/omp_get_initial_device(), Size, Code);
}
}
@@ -225,7 +225,7 @@ void Interface::endTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin,
ompt_scope_end, TargetTaskData, &TargetData, &HostOpId,
ompt_target_data_transfer_from_device, TgtPtrBegin, DeviceId,
HstPtrBegin,
- /* TgtDeviceNum */ omp_get_initial_device(), Size, Code);
+ /*TgtDeviceNum=*/omp_get_initial_device(), Size, Code);
}
endTargetDataOperation();
}
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index fa8932361a51c9..654efd524024a4 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -117,7 +117,7 @@ void *DeviceTy::allocData(int64_t Size, void *HstPtr, int32_t Kind) {
OMPT_IF_BUILT(InterfaceRAII TargetDataAllocRAII(
RegionInterface.getCallbacks<ompt_target_data_alloc>(),
DeviceID, HstPtr, &TargetPtr, Size,
- /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
+ /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));)
TargetPtr = RTL->data_alloc(RTLDeviceID, Size, HstPtr, Kind);
return TargetPtr;
@@ -128,7 +128,7 @@ int32_t DeviceTy::deleteData(void *TgtAllocBegin, int32_t Kind) {
OMPT_IF_BUILT(InterfaceRAII TargetDataDeleteRAII(
RegionInterface.getCallbacks<ompt_target_data_delete>(),
DeviceID, TgtAllocBegin,
- /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
+ /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));)
return RTL->data_delete(RTLDeviceID, TgtAllocBegin, Kind);
}
@@ -146,7 +146,7 @@ int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
InterfaceRAII TargetDataSubmitRAII(
RegionInterface.getCallbacks<ompt_target_data_transfer_to_device>(),
DeviceID, TgtPtrBegin, HstPtrBegin, Size,
- /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
+ /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));)
if (!AsyncInfo || !RTL->data_submit_async || !RTL->synchronize)
return RTL->data_submit(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size);
@@ -168,7 +168,7 @@ int32_t DeviceTy::retrieveData(void *HstPtrBegin, void *TgtPtrBegin,
InterfaceRAII TargetDataRetrieveRAII(
RegionInterface.getCallbacks<ompt_target_data_transfer_from_device>(),
DeviceID, HstPtrBegin, TgtPtrBegin, Size,
- /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
+ /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));)
if (!RTL->data_retrieve_async || !RTL->synchronize)
return RTL->data_retrieve(RTLDeviceID, HstPtrBegin, TgtPtrBegin, Size);
diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp
index 61d9db17f51006..49495ac266f1b3 100644
--- a/openmp/libomptarget/src/interface.cpp
+++ b/openmp/libomptarget/src/interface.cpp
@@ -113,7 +113,7 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
int Rc = OFFLOAD_SUCCESS;
Rc = TargetDataFunction(Loc, *DeviceOrErr, ArgNum, ArgsBase, Args, ArgSizes,
ArgTypes, ArgNames, ArgMappers, AsyncInfo,
- false /* FromMapper */);
+ false /*FromMapper=*/);
if (Rc == OFFLOAD_SUCCESS)
Rc = AsyncInfo.synchronize();
@@ -293,7 +293,7 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
/// RAII to establish tool anchors before and after target region
OMPT_IF_BUILT(InterfaceRAII TargetRAII(
RegionInterface.getCallbacks<ompt_target>(), DeviceId,
- /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
+ /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));)
int Rc = OFFLOAD_SUCCESS;
Rc = target(Loc, *DeviceOrErr, HostPtr, *KernelArgs, AsyncInfo);
@@ -387,7 +387,7 @@ EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId,
/// RAII to establish tool anchors before and after target region
OMPT_IF_BUILT(InterfaceRAII TargetRAII(
RegionInterface.getCallbacks<ompt_target>(), DeviceId,
- /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
+ /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));)
AsyncInfoTy AsyncInfo(*DeviceOrErr);
int Rc = target_replay(Loc, *DeviceOrErr, HostPtr, DeviceMemory,
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index a7d55d7ebd5391..ed34870e12d011 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -225,7 +225,7 @@ static int initLibrary(DeviceTy &Device) {
AsyncInfoTy AsyncInfo(Device);
void *DevPtr;
Device.retrieveData(&DevPtr, CurrDeviceEntryAddr, sizeof(void *),
- AsyncInfo, /* Entry */ nullptr, &HDTTMap);
+ AsyncInfo, /*Entry=*/nullptr, &HDTTMap);
if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
CurrDeviceEntryAddr = DevPtr;
@@ -617,7 +617,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
/*HstPtrName=*/nullptr,
/*HasFlagTo=*/false, /*HasFlagAlways=*/false, IsImplicit, UpdateRef,
HasCloseModifier, HasPresentModifier, HasHoldModifier, AsyncInfo,
- /* OwnedTPR */ nullptr, /* ReleaseHDTTMap */ false);
+ /*OwnedTPR=*/nullptr, /*ReleaseHDTTMap=*/false);
PointerTgtPtrBegin = PointerTpr.TargetPointer;
IsHostPtr = PointerTpr.Flags.IsHostPointer;
if (!PointerTgtPtrBegin) {
@@ -1723,7 +1723,7 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
TargetTable->EntriesBegin[TM->Index].name, DPxPTR(TgtEntryPtr), TM->Index);
- void *TgtPtr = Device.allocData(DeviceMemorySize, /* HstPtr */ nullptr,
+ void *TgtPtr = Device.allocData(DeviceMemorySize, /*HstPtr=*/nullptr,
TARGET_ALLOC_DEFAULT);
Device.submitData(TgtPtr, DeviceMemory, DeviceMemorySize, AsyncInfo);
diff --git a/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp b/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp
index fc3ff078a81fb9..761e04e4c7bbdb 100644
--- a/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp
+++ b/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp
@@ -57,8 +57,8 @@ int main(int argc, char **argv) {
cl::ParseCommandLineOptions(argc, argv, "llvm-omp-kernel-replay\n");
ErrorOr<std::unique_ptr<MemoryBuffer>> KernelInfoMB =
- MemoryBuffer::getFile(InputFilename, /* isText */ true,
- /* RequiresNullTerminator */ true);
+ MemoryBuffer::getFile(InputFilename, /*isText=*/true,
+ /*RequiresNullTerminator=*/true);
if (!KernelInfoMB)
report_fatal_error("Error reading the kernel info json file");
Expected<json::Value> JsonKernelInfo =
@@ -100,8 +100,8 @@ int main(int argc, char **argv) {
KernelEntry.addr = (void *)0x1;
ErrorOr<std::unique_ptr<MemoryBuffer>> ImageMB =
- MemoryBuffer::getFile(KernelEntryName + ".image", /* isText */ false,
- /* RequiresNullTerminator */ false);
+ MemoryBuffer::getFile(KernelEntryName + ".image", /*isText=*/false,
+ /*RequiresNullTerminator=*/false);
if (!ImageMB)
report_fatal_error("Error reading the kernel image.");
@@ -127,7 +127,7 @@ int main(int argc, char **argv) {
int32_t DeviceId = (DeviceIdOpt > -1 ? DeviceIdOpt : DeviceIdJson.value());
// TODO: do we need requires?
- //__tgt_register_requires(/* Flags */1);
+ //__tgt_register_requires(/*Flags=*/1);
__tgt_register_lib(&Desc);
@@ -140,8 +140,8 @@ int main(int argc, char **argv) {
}
ErrorOr<std::unique_ptr<MemoryBuffer>> DeviceMemoryMB =
- MemoryBuffer::getFile(KernelEntryName + ".memory", /* isText */ false,
- /* RequiresNullTerminator */ false);
+ MemoryBuffer::getFile(KernelEntryName + ".memory", /*isText=*/false,
+ /*RequiresNullTerminator=*/false);
if (!DeviceMemoryMB)
report_fatal_error("Error reading the kernel input device memory.");
@@ -166,7 +166,7 @@ int main(int argc, char **argv) {
}
__tgt_target_kernel_replay(
- /* Loc */ nullptr, DeviceId, KernelEntry.addr, (char *)recored_data,
+ /*Loc=*/nullptr, DeviceId, KernelEntry.addr, (char *)recored_data,
DeviceMemoryMB.get()->getBufferSize(), TgtArgs.data(),
TgtArgOffsets.data(), NumArgs.value(), NumTeams, NumThreads,
LoopTripCount.value());
@@ -174,15 +174,15 @@ int main(int argc, char **argv) {
if (VerifyOpt) {
ErrorOr<std::unique_ptr<MemoryBuffer>> OriginalOutputMB =
MemoryBuffer::getFile(KernelEntryName + ".original.output",
- /* isText */ false,
- /* RequiresNullTerminator */ false);
+ /*isText=*/false,
+ /*RequiresNullTerminator=*/false);
if (!OriginalOutputMB)
report_fatal_error("Error reading the kernel original output file, make "
"sure LIBOMPTARGET_SAVE_OUTPUT is set when recording");
ErrorOr<std::unique_ptr<MemoryBuffer>> ReplayOutputMB =
MemoryBuffer::getFile(KernelEntryName + ".replay.output",
- /* isText */ false,
- /* RequiresNullTerminator */ false);
+ /*isText=*/false,
+ /*RequiresNullTerminator=*/false);
if (!ReplayOutputMB)
report_fatal_error("Error reading the kernel replay output file");
More information about the Openmp-commits
mailing list