[Openmp-commits] [openmp] 46642cc - [Libomptarget] Remove debug RAII from libomptarget
Joseph Huber via Openmp-commits
openmp-commits at lists.llvm.org
Thu Aug 3 07:37:54 PDT 2023
Author: Joseph Huber
Date: 2023-08-03T09:37:47-05:00
New Revision: 46642cc83dc575962e1a6eb557714319c65ca5b8
URL: https://github.com/llvm/llvm-project/commit/46642cc83dc575962e1a6eb557714319c65ca5b8
DIFF: https://github.com/llvm/llvm-project/commit/46642cc83dc575962e1a6eb557714319c65ca5b8.diff
LOG: [Libomptarget] Remove debug RAII from libomptarget
This feature was supposed to allow you to trace execution inside of
Libomptarget. However, this never really worked properly. The printing
was always reoganized, only worked for single threads, and pretty much
only told you a handful of things about a runtime library that's an
implementation detail to all users. Despite this, it contributed about
40% of the total filesize of the deviceRTL. This patch simply removes
this functionalit which I think was past due.
Reviewed By: tianshilei1992
Differential Revision: https://reviews.llvm.org/D157001
Added:
Modified:
openmp/docs/design/Runtimes.rst
openmp/libomptarget/DeviceRTL/include/Debug.h
openmp/libomptarget/DeviceRTL/src/Debug.cpp
openmp/libomptarget/DeviceRTL/src/Kernel.cpp
openmp/libomptarget/DeviceRTL/src/Mapping.cpp
openmp/libomptarget/DeviceRTL/src/Misc.cpp
openmp/libomptarget/DeviceRTL/src/Parallelism.cpp
openmp/libomptarget/DeviceRTL/src/Reduction.cpp
openmp/libomptarget/DeviceRTL/src/State.cpp
openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
openmp/libomptarget/DeviceRTL/src/Tasking.cpp
openmp/libomptarget/DeviceRTL/src/Utils.cpp
openmp/libomptarget/DeviceRTL/src/Workshare.cpp
Removed:
################################################################################
diff --git a/openmp/docs/design/Runtimes.rst b/openmp/docs/design/Runtimes.rst
index 84a65b183bf046..e681252731dcbd 100644
--- a/openmp/docs/design/Runtimes.rst
+++ b/openmp/docs/design/Runtimes.rst
@@ -1452,34 +1452,4 @@ to selectively enable and disable
diff erent features. Currently, the following
debugging features are supported.
* Enable debugging assertions in the device. ``0x01``
- * Enable OpenMP runtime function traces in the device. ``0x2``
* Enable diagnosing common problems during offloading . ``0x4``
-
-.. code-block:: c++
-
- void copy(double *X, double *Y) {
- #pragma omp target teams distribute parallel for
- for (std::size_t i = 0; i < N; ++i)
- Y[i] = X[i];
- }
-
-Compiling this code targeting ``nvptx64`` with debugging enabled will
-provide the following output from the device runtime library.
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -fopenmp-target-debug=3
- $ env LIBOMPTARGET_DEVICE_RTL_DEBUG=3 ./zaxpy
-
-.. code-block:: text
-
- Kernel.cpp:70: Thread 0 Entering int32_t __kmpc_target_init()
- Parallelism.cpp:196: Thread 0 Entering int32_t __kmpc_global_thread_num()
- Mapping.cpp:239: Thread 0 Entering uint32_t __kmpc_get_hardware_num_threads_in_block()
- Workshare.cpp:616: Thread 0 Entering void __kmpc_distribute_static_init_4()
- Parallelism.cpp:85: Thread 0 Entering void __kmpc_parallel_51()
- Parallelism.cpp:69: Thread 0 Entering <OpenMP Outlined Function>
- Workshare.cpp:575: Thread 0 Entering void __kmpc_for_static_init_4()
- Workshare.cpp:660: Thread 0 Entering void __kmpc_distribute_static_fini()
- Workshare.cpp:660: Thread 0 Entering void __kmpc_distribute_static_fini()
- Kernel.cpp:103: Thread 0 Entering void __kmpc_target_deinit()
diff --git a/openmp/libomptarget/DeviceRTL/include/Debug.h b/openmp/libomptarget/DeviceRTL/include/Debug.h
index 1ea129a97a17e8..54628c03a3cc55 100644
--- a/openmp/libomptarget/DeviceRTL/include/Debug.h
+++ b/openmp/libomptarget/DeviceRTL/include/Debug.h
@@ -42,17 +42,4 @@ void __assert_fail(const char *expr, const char *msg, const char *file,
///}
-/// Enter a debugging scope for performing function traces. Enabled with
-/// FunctionTracting set in the debug kind.
-#define FunctionTracingRAII() \
- DebugEntryRAII Entry(__FILE__, __LINE__, __PRETTY_FUNCTION__);
-
-/// An RAII class for handling entries to debug locations. The current location
-/// and function will be printed on entry. Nested levels increase the
-/// indentation shown in the debugging output.
-struct DebugEntryRAII {
- DebugEntryRAII(const char *File, const unsigned Line, const char *Function);
- ~DebugEntryRAII();
-};
-
#endif
diff --git a/openmp/libomptarget/DeviceRTL/src/Debug.cpp b/openmp/libomptarget/DeviceRTL/src/Debug.cpp
index 6e296b0a277cae..4ed5d742165ec1 100644
--- a/openmp/libomptarget/DeviceRTL/src/Debug.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Debug.cpp
@@ -37,32 +37,4 @@ void __assert_fail(const char *expr, const char *msg, const char *file,
}
}
-DebugEntryRAII::DebugEntryRAII(const char *File, const unsigned Line,
- const char *Function) {
- if (config::isDebugMode(config::DebugKind::FunctionTracing) &&
- mapping::getThreadIdInBlock() == 0 &&
- mapping::getBlockIdInKernel() == 0) {
-
- uint16_t &Level =
- state::getKernelEnvironment().DynamicEnv->DebugIndentionLevel;
-
- for (int I = 0; I < Level; ++I)
- PRINTF("%s", " ");
-
- PRINTF("%s:%u: Thread %u Entering %s\n", File, Line,
- mapping::getThreadIdInBlock(), Function);
- Level++;
- }
-}
-
-DebugEntryRAII::~DebugEntryRAII() {
- if (config::isDebugMode(config::DebugKind::FunctionTracing) &&
- mapping::getThreadIdInBlock() == 0 &&
- mapping::getBlockIdInKernel() == 0) {
- uint16_t &Level =
- state::getKernelEnvironment().DynamicEnv->DebugIndentionLevel;
- Level--;
- }
-}
-
#pragma omp end declare target
diff --git a/openmp/libomptarget/DeviceRTL/src/Kernel.cpp b/openmp/libomptarget/DeviceRTL/src/Kernel.cpp
index e8f6cfdc6ea540..93a09972af1662 100644
--- a/openmp/libomptarget/DeviceRTL/src/Kernel.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Kernel.cpp
@@ -34,8 +34,6 @@ static void inititializeRuntime(bool IsSPMD,
/// Simple generic state machine for worker threads.
static void genericStateMachine(IdentTy *Ident) {
- FunctionTracingRAII();
-
uint32_t TId = mapping::getThreadIdInBlock();
do {
@@ -70,7 +68,6 @@ extern "C" {
/// \param Ident Source location identification, can be NULL.
///
int32_t __kmpc_target_init(KernelEnvironmentTy &KernelEnvironment) {
- FunctionTracingRAII();
ConfigurationEnvironmentTy &Configuration = KernelEnvironment.Configuration;
bool IsSPMD = Configuration.ExecMode &
llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD;
@@ -137,7 +134,6 @@ int32_t __kmpc_target_init(KernelEnvironmentTy &KernelEnvironment) {
/// \param Ident Source location identification, can be NULL.
///
void __kmpc_target_deinit() {
- FunctionTracingRAII();
bool IsSPMD = mapping::isSPMDMode();
state::assumeInitialState(IsSPMD);
if (IsSPMD)
@@ -147,10 +143,7 @@ void __kmpc_target_deinit() {
state::ParallelRegionFn = nullptr;
}
-int8_t __kmpc_is_spmd_exec_mode() {
- FunctionTracingRAII();
- return mapping::isSPMDMode();
-}
+int8_t __kmpc_is_spmd_exec_mode() { return mapping::isSPMDMode(); }
}
#pragma omp end declare target
diff --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
index 2f50530e79a1d0..cf1689e4ccd138 100644
--- a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
@@ -357,17 +357,14 @@ bool mapping::isGenericMode() { return !isSPMDMode(); }
extern "C" {
__attribute__((noinline)) uint32_t __kmpc_get_hardware_thread_id_in_block() {
- FunctionTracingRAII();
return mapping::getThreadIdInBlock();
}
__attribute__((noinline)) uint32_t __kmpc_get_hardware_num_threads_in_block() {
- FunctionTracingRAII();
return impl::getNumberOfThreadsInBlock(mapping::DIM_X);
}
__attribute__((noinline)) uint32_t __kmpc_get_warp_size() {
- FunctionTracingRAII();
return impl::getWarpSize();
}
}
diff --git a/openmp/libomptarget/DeviceRTL/src/Misc.cpp b/openmp/libomptarget/DeviceRTL/src/Misc.cpp
index a19a263e55b246..0c361fe61061de 100644
--- a/openmp/libomptarget/DeviceRTL/src/Misc.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Misc.cpp
@@ -77,15 +77,9 @@ double getWTime() {
///{
extern "C" {
-int32_t __kmpc_cancellationpoint(IdentTy *, int32_t, int32_t) {
- FunctionTracingRAII();
- return 0;
-}
+int32_t __kmpc_cancellationpoint(IdentTy *, int32_t, int32_t) { return 0; }
-int32_t __kmpc_cancel(IdentTy *, int32_t, int32_t) {
- FunctionTracingRAII();
- return 0;
-}
+int32_t __kmpc_cancel(IdentTy *, int32_t, int32_t) { return 0; }
double omp_get_wtick(void) { return ompx::impl::getWTick(); }
diff --git a/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp b/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp
index 84d8f88105f107..9ff807038a99c4 100644
--- a/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp
@@ -70,7 +70,6 @@ uint32_t determineNumberOfThreads(int32_t NumThreadsClause) {
// Invoke an outlined parallel function unwrapping arguments (up to 32).
void invokeMicrotask(int32_t global_tid, int32_t bound_tid, void *fn,
void **args, int64_t nargs) {
- DebugEntryRAII Entry(__FILE__, __LINE__, "<OpenMP Outlined Function>");
switch (nargs) {
#include "generated_microtask_cases.gen"
default:
@@ -86,8 +85,6 @@ extern "C" {
void __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr,
int32_t num_threads, int proc_bind, void *fn,
void *wrapper_fn, void **args, int64_t nargs) {
- FunctionTracingRAII();
-
uint32_t TId = mapping::getThreadIdInBlock();
// Assert the parallelism level is zero if disabled by the user.
@@ -264,7 +261,6 @@ void __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr,
__attribute__((noinline)) bool
__kmpc_kernel_parallel(ParallelRegionFnTy *WorkFn) {
- FunctionTracingRAII();
// Work function and arguments for L1 parallel region.
*WorkFn = state::ParallelRegionFn;
@@ -279,7 +275,6 @@ __kmpc_kernel_parallel(ParallelRegionFnTy *WorkFn) {
}
__attribute__((noinline)) void __kmpc_kernel_end_parallel() {
- FunctionTracingRAII();
// In case we have modified an ICV for this thread before a ThreadState was
// created. We drop it now to not contaminate the next parallel region.
ASSERT(!mapping::isSPMDMode(), nullptr);
@@ -288,24 +283,14 @@ __attribute__((noinline)) void __kmpc_kernel_end_parallel() {
ASSERT(!mapping::isSPMDMode(), nullptr);
}
-uint16_t __kmpc_parallel_level(IdentTy *, uint32_t) {
- FunctionTracingRAII();
- return omp_get_level();
-}
+uint16_t __kmpc_parallel_level(IdentTy *, uint32_t) { return omp_get_level(); }
-int32_t __kmpc_global_thread_num(IdentTy *) {
- FunctionTracingRAII();
- return omp_get_thread_num();
-}
+int32_t __kmpc_global_thread_num(IdentTy *) { return omp_get_thread_num(); }
void __kmpc_push_num_teams(IdentTy *loc, int32_t tid, int32_t num_teams,
- int32_t thread_limit) {
- FunctionTracingRAII();
-}
+ int32_t thread_limit) {}
-void __kmpc_push_proc_bind(IdentTy *loc, uint32_t tid, int proc_bind) {
- FunctionTracingRAII();
-}
+void __kmpc_push_proc_bind(IdentTy *loc, uint32_t tid, int proc_bind) {}
}
#pragma omp end declare target
diff --git a/openmp/libomptarget/DeviceRTL/src/Reduction.cpp b/openmp/libomptarget/DeviceRTL/src/Reduction.cpp
index 8e0b91bfc1eb14..8ec8d296196337 100644
--- a/openmp/libomptarget/DeviceRTL/src/Reduction.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Reduction.cpp
@@ -176,7 +176,6 @@ extern "C" {
int32_t __kmpc_nvptx_parallel_reduce_nowait_v2(
IdentTy *Loc, int32_t TId, int32_t num_vars, uint64_t reduce_size,
void *reduce_data, ShuffleReductFnTy shflFct, InterWarpCopyFnTy cpyFct) {
- FunctionTracingRAII();
return nvptx_parallel_reduce_nowait(TId, num_vars, reduce_size, reduce_data,
shflFct, cpyFct, mapping::isSPMDMode(),
false);
@@ -187,8 +186,6 @@ int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
void *reduce_data, ShuffleReductFnTy shflFct, InterWarpCopyFnTy cpyFct,
ListGlobalFnTy lgcpyFct, ListGlobalFnTy lgredFct, ListGlobalFnTy glcpyFct,
ListGlobalFnTy glredFct) {
- FunctionTracingRAII();
-
// Terminate all threads in non-SPMD mode except for the master thread.
uint32_t ThreadId = mapping::getThreadIdInBlock();
if (mapping::isGenericMode()) {
@@ -311,9 +308,9 @@ int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
return 0;
}
-void __kmpc_nvptx_end_reduce(int32_t TId) { FunctionTracingRAII(); }
+void __kmpc_nvptx_end_reduce(int32_t TId) {}
-void __kmpc_nvptx_end_reduce_nowait(int32_t TId) { FunctionTracingRAII(); }
+void __kmpc_nvptx_end_reduce_nowait(int32_t TId) {}
}
#pragma omp end declare target
diff --git a/openmp/libomptarget/DeviceRTL/src/State.cpp b/openmp/libomptarget/DeviceRTL/src/State.cpp
index c18368da72d6cd..721137cb95d658 100644
--- a/openmp/libomptarget/DeviceRTL/src/State.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/State.cpp
@@ -424,12 +424,10 @@ int omp_get_initial_device(void) { return -1; }
extern "C" {
__attribute__((noinline)) void *__kmpc_alloc_shared(uint64_t Bytes) {
- FunctionTracingRAII();
return memory::allocShared(Bytes, "Frontend alloc shared");
}
__attribute__((noinline)) void __kmpc_free_shared(void *Ptr, uint64_t Bytes) {
- FunctionTracingRAII();
memory::freeShared(Ptr, Bytes, "Frontend free shared");
}
@@ -455,7 +453,6 @@ constexpr uint64_t NUM_SHARED_VARIABLES_IN_SHARED_MEM = 64;
allocator(omp_pteam_mem_alloc)
void __kmpc_begin_sharing_variables(void ***GlobalArgs, uint64_t nArgs) {
- FunctionTracingRAII();
if (nArgs <= NUM_SHARED_VARIABLES_IN_SHARED_MEM) {
SharedMemVariableSharingSpacePtr = &SharedMemVariableSharingSpace[0];
} else {
@@ -468,13 +465,11 @@ void __kmpc_begin_sharing_variables(void ***GlobalArgs, uint64_t nArgs) {
}
void __kmpc_end_sharing_variables() {
- FunctionTracingRAII();
if (SharedMemVariableSharingSpacePtr != &SharedMemVariableSharingSpace[0])
memory::freeGlobal(SharedMemVariableSharingSpacePtr, "new extended args");
}
void __kmpc_get_shared_variables(void ***GlobalArgs) {
- FunctionTracingRAII();
*GlobalArgs = SharedMemVariableSharingSpacePtr;
}
}
diff --git a/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp b/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
index 695f77bab215cd..3370c5a8472f0b 100644
--- a/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
@@ -504,18 +504,16 @@ void unsetCriticalLock(omp_lock_t *Lock) { impl::unsetLock(Lock); }
void setCriticalLock(omp_lock_t *Lock) { impl::setLock(Lock); }
extern "C" {
-void __kmpc_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
+void __kmpc_ordered(IdentTy *Loc, int32_t TId) {}
-void __kmpc_end_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
+void __kmpc_end_ordered(IdentTy *Loc, int32_t TId) {}
int32_t __kmpc_cancel_barrier(IdentTy *Loc, int32_t TId) {
- FunctionTracingRAII();
__kmpc_barrier(Loc, TId);
return 0;
}
void __kmpc_barrier(IdentTy *Loc, int32_t TId) {
- FunctionTracingRAII();
if (mapping::isMainThreadInGenericMode())
return __kmpc_flush(Loc);
@@ -527,62 +525,45 @@ void __kmpc_barrier(IdentTy *Loc, int32_t TId) {
__attribute__((noinline)) void __kmpc_barrier_simple_spmd(IdentTy *Loc,
int32_t TId) {
- FunctionTracingRAII();
synchronize::threadsAligned(atomic::OrderingTy::seq_cst);
}
__attribute__((noinline)) void __kmpc_barrier_simple_generic(IdentTy *Loc,
int32_t TId) {
- FunctionTracingRAII();
synchronize::threads(atomic::OrderingTy::seq_cst);
}
int32_t __kmpc_master(IdentTy *Loc, int32_t TId) {
- FunctionTracingRAII();
return omp_get_thread_num() == 0;
}
-void __kmpc_end_master(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
+void __kmpc_end_master(IdentTy *Loc, int32_t TId) {}
int32_t __kmpc_masked(IdentTy *Loc, int32_t TId, int32_t Filter) {
- FunctionTracingRAII();
return omp_get_thread_num() == Filter;
}
-void __kmpc_end_masked(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
+void __kmpc_end_masked(IdentTy *Loc, int32_t TId) {}
int32_t __kmpc_single(IdentTy *Loc, int32_t TId) {
- FunctionTracingRAII();
return __kmpc_master(Loc, TId);
}
void __kmpc_end_single(IdentTy *Loc, int32_t TId) {
- FunctionTracingRAII();
// The barrier is explicitly called.
}
-void __kmpc_flush(IdentTy *Loc) {
- FunctionTracingRAII();
- fence::kernel(atomic::seq_cst);
-}
+void __kmpc_flush(IdentTy *Loc) { fence::kernel(atomic::seq_cst); }
-uint64_t __kmpc_warp_active_thread_mask(void) {
- FunctionTracingRAII();
- return mapping::activemask();
-}
+uint64_t __kmpc_warp_active_thread_mask(void) { return mapping::activemask(); }
-void __kmpc_syncwarp(uint64_t Mask) {
- FunctionTracingRAII();
- synchronize::warp(Mask);
-}
+void __kmpc_syncwarp(uint64_t Mask) { synchronize::warp(Mask); }
void __kmpc_critical(IdentTy *Loc, int32_t TId, CriticalNameTy *Name) {
- FunctionTracingRAII();
impl::setCriticalLock(reinterpret_cast<omp_lock_t *>(Name));
}
void __kmpc_end_critical(IdentTy *Loc, int32_t TId, CriticalNameTy *Name) {
- FunctionTracingRAII();
impl::unsetCriticalLock(reinterpret_cast<omp_lock_t *>(Name));
}
diff --git a/openmp/libomptarget/DeviceRTL/src/Tasking.cpp b/openmp/libomptarget/DeviceRTL/src/Tasking.cpp
index 6555eae7fc2808..2dc33562e6d796 100644
--- a/openmp/libomptarget/DeviceRTL/src/Tasking.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Tasking.cpp
@@ -28,7 +28,6 @@ TaskDescriptorTy *__kmpc_omp_task_alloc(IdentTy *, int32_t, int32_t,
size_t TaskSizeInclPrivateValues,
size_t SharedValuesSize,
TaskFnTy TaskFn) {
- FunctionTracingRAII();
auto TaskSizeInclPrivateValuesPadded =
utils::roundUp(TaskSizeInclPrivateValues, uint64_t(sizeof(void *)));
auto TaskSizeTotal = TaskSizeInclPrivateValuesPadded + SharedValuesSize;
@@ -43,14 +42,12 @@ TaskDescriptorTy *__kmpc_omp_task_alloc(IdentTy *, int32_t, int32_t,
int32_t __kmpc_omp_task(IdentTy *Loc, uint32_t TId,
TaskDescriptorTy *TaskDescriptor) {
- FunctionTracingRAII();
return __kmpc_omp_task_with_deps(Loc, TId, TaskDescriptor, 0, 0, 0, 0);
}
int32_t __kmpc_omp_task_with_deps(IdentTy *Loc, uint32_t TId,
TaskDescriptorTy *TaskDescriptor, int32_t,
void *, int32_t, void *) {
- FunctionTracingRAII();
state::DateEnvironmentRAII DERAII(Loc);
TaskDescriptor->TaskFn(0, TaskDescriptor);
@@ -61,42 +58,31 @@ int32_t __kmpc_omp_task_with_deps(IdentTy *Loc, uint32_t TId,
void __kmpc_omp_task_begin_if0(IdentTy *Loc, uint32_t TId,
TaskDescriptorTy *TaskDescriptor) {
- FunctionTracingRAII();
state::enterDataEnvironment(Loc);
}
void __kmpc_omp_task_complete_if0(IdentTy *Loc, uint32_t TId,
TaskDescriptorTy *TaskDescriptor) {
- FunctionTracingRAII();
state::exitDataEnvironment();
memory::freeGlobal(TaskDescriptor, "explicit task descriptor");
}
void __kmpc_omp_wait_deps(IdentTy *Loc, uint32_t TId, int32_t, void *, int32_t,
- void *) {
- FunctionTracingRAII();
-}
+ void *) {}
-void __kmpc_taskgroup(IdentTy *Loc, uint32_t TId) { FunctionTracingRAII(); }
+void __kmpc_taskgroup(IdentTy *Loc, uint32_t TId) {}
-void __kmpc_end_taskgroup(IdentTy *Loc, uint32_t TId) { FunctionTracingRAII(); }
+void __kmpc_end_taskgroup(IdentTy *Loc, uint32_t TId) {}
-int32_t __kmpc_omp_taskyield(IdentTy *Loc, uint32_t TId, int) {
- FunctionTracingRAII();
- return 0;
-}
+int32_t __kmpc_omp_taskyield(IdentTy *Loc, uint32_t TId, int) { return 0; }
-int32_t __kmpc_omp_taskwait(IdentTy *Loc, uint32_t TId) {
- FunctionTracingRAII();
- return 0;
-}
+int32_t __kmpc_omp_taskwait(IdentTy *Loc, uint32_t TId) { return 0; }
void __kmpc_taskloop(IdentTy *Loc, uint32_t TId,
TaskDescriptorTy *TaskDescriptor, int,
uint64_t *LowerBound, uint64_t *UpperBound, int64_t, int,
int32_t, uint64_t, void *) {
- FunctionTracingRAII();
// Skip task entirely if empty iteration space.
if (*LowerBound > *UpperBound)
return;
diff --git a/openmp/libomptarget/DeviceRTL/src/Utils.cpp b/openmp/libomptarget/DeviceRTL/src/Utils.cpp
index d74f7e069cf67c..6125236863098f 100644
--- a/openmp/libomptarget/DeviceRTL/src/Utils.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Utils.cpp
@@ -144,12 +144,10 @@ bool utils::isSharedMemPtr(void *Ptr) { return impl::isSharedMemPtr(Ptr); }
extern "C" {
int32_t __kmpc_shuffle_int32(int32_t Val, int16_t Delta, int16_t SrcLane) {
- FunctionTracingRAII();
return impl::shuffleDown(lanes::All, Val, Delta, SrcLane);
}
int64_t __kmpc_shuffle_int64(int64_t Val, int16_t Delta, int16_t Width) {
- FunctionTracingRAII();
uint32_t lo, hi;
utils::unpack(Val, lo, hi);
hi = impl::shuffleDown(lanes::All, hi, Delta, Width);
diff --git a/openmp/libomptarget/DeviceRTL/src/Workshare.cpp b/openmp/libomptarget/DeviceRTL/src/Workshare.cpp
index c56c59d4b3ed4a..0dbfafc4d699e7 100644
--- a/openmp/libomptarget/DeviceRTL/src/Workshare.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Workshare.cpp
@@ -473,7 +473,6 @@ extern "C" {
// init
void __kmpc_dispatch_init_4(IdentTy *loc, int32_t tid, int32_t schedule,
int32_t lb, int32_t ub, int32_t st, int32_t chunk) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = pushDST();
omptarget_nvptx_LoopSupport<int32_t, int32_t>::dispatch_init(
loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk, DST);
@@ -482,7 +481,6 @@ void __kmpc_dispatch_init_4(IdentTy *loc, int32_t tid, int32_t schedule,
void __kmpc_dispatch_init_4u(IdentTy *loc, int32_t tid, int32_t schedule,
uint32_t lb, uint32_t ub, int32_t st,
int32_t chunk) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = pushDST();
omptarget_nvptx_LoopSupport<uint32_t, int32_t>::dispatch_init(
loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk, DST);
@@ -490,7 +488,6 @@ void __kmpc_dispatch_init_4u(IdentTy *loc, int32_t tid, int32_t schedule,
void __kmpc_dispatch_init_8(IdentTy *loc, int32_t tid, int32_t schedule,
int64_t lb, int64_t ub, int64_t st, int64_t chunk) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = pushDST();
omptarget_nvptx_LoopSupport<int64_t, int64_t>::dispatch_init(
loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk, DST);
@@ -499,7 +496,6 @@ void __kmpc_dispatch_init_8(IdentTy *loc, int32_t tid, int32_t schedule,
void __kmpc_dispatch_init_8u(IdentTy *loc, int32_t tid, int32_t schedule,
uint64_t lb, uint64_t ub, int64_t st,
int64_t chunk) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = pushDST();
omptarget_nvptx_LoopSupport<uint64_t, int64_t>::dispatch_init(
loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk, DST);
@@ -508,7 +504,6 @@ void __kmpc_dispatch_init_8u(IdentTy *loc, int32_t tid, int32_t schedule,
// next
int __kmpc_dispatch_next_4(IdentTy *loc, int32_t tid, int32_t *p_last,
int32_t *p_lb, int32_t *p_ub, int32_t *p_st) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = peekDST();
return omptarget_nvptx_LoopSupport<int32_t, int32_t>::dispatch_next(
loc, tid, p_last, p_lb, p_ub, p_st, DST);
@@ -516,7 +511,6 @@ int __kmpc_dispatch_next_4(IdentTy *loc, int32_t tid, int32_t *p_last,
int __kmpc_dispatch_next_4u(IdentTy *loc, int32_t tid, int32_t *p_last,
uint32_t *p_lb, uint32_t *p_ub, int32_t *p_st) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = peekDST();
return omptarget_nvptx_LoopSupport<uint32_t, int32_t>::dispatch_next(
loc, tid, p_last, p_lb, p_ub, p_st, DST);
@@ -524,7 +518,6 @@ int __kmpc_dispatch_next_4u(IdentTy *loc, int32_t tid, int32_t *p_last,
int __kmpc_dispatch_next_8(IdentTy *loc, int32_t tid, int32_t *p_last,
int64_t *p_lb, int64_t *p_ub, int64_t *p_st) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = peekDST();
return omptarget_nvptx_LoopSupport<int64_t, int64_t>::dispatch_next(
loc, tid, p_last, p_lb, p_ub, p_st, DST);
@@ -532,7 +525,6 @@ int __kmpc_dispatch_next_8(IdentTy *loc, int32_t tid, int32_t *p_last,
int __kmpc_dispatch_next_8u(IdentTy *loc, int32_t tid, int32_t *p_last,
uint64_t *p_lb, uint64_t *p_ub, int64_t *p_st) {
- FunctionTracingRAII();
DynamicScheduleTracker *DST = peekDST();
return omptarget_nvptx_LoopSupport<uint64_t, int64_t>::dispatch_next(
loc, tid, p_last, p_lb, p_ub, p_st, DST);
@@ -540,25 +532,21 @@ int __kmpc_dispatch_next_8u(IdentTy *loc, int32_t tid, int32_t *p_last,
// fini
void __kmpc_dispatch_fini_4(IdentTy *loc, int32_t tid) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<int32_t, int32_t>::dispatch_fini();
popDST();
}
void __kmpc_dispatch_fini_4u(IdentTy *loc, int32_t tid) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<uint32_t, int32_t>::dispatch_fini();
popDST();
}
void __kmpc_dispatch_fini_8(IdentTy *loc, int32_t tid) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<int64_t, int64_t>::dispatch_fini();
popDST();
}
void __kmpc_dispatch_fini_8u(IdentTy *loc, int32_t tid) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<uint64_t, int64_t>::dispatch_fini();
popDST();
}
@@ -571,7 +559,6 @@ void __kmpc_for_static_init_4(IdentTy *loc, int32_t global_tid,
int32_t schedtype, int32_t *plastiter,
int32_t *plower, int32_t *pupper,
int32_t *pstride, int32_t incr, int32_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
@@ -581,7 +568,6 @@ void __kmpc_for_static_init_4u(IdentTy *loc, int32_t global_tid,
int32_t schedtype, int32_t *plastiter,
uint32_t *plower, uint32_t *pupper,
int32_t *pstride, int32_t incr, int32_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
@@ -591,7 +577,6 @@ void __kmpc_for_static_init_8(IdentTy *loc, int32_t global_tid,
int32_t schedtype, int32_t *plastiter,
int64_t *plower, int64_t *pupper,
int64_t *pstride, int64_t incr, int64_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
@@ -601,7 +586,6 @@ void __kmpc_for_static_init_8u(IdentTy *loc, int32_t global_tid,
int32_t schedtype, int32_t *plastiter,
uint64_t *plower, uint64_t *pupper,
int64_t *pstride, int64_t incr, int64_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
@@ -612,7 +596,6 @@ void __kmpc_distribute_static_init_4(IdentTy *loc, int32_t global_tid,
int32_t *plower, int32_t *pupper,
int32_t *pstride, int32_t incr,
int32_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
@@ -623,7 +606,6 @@ void __kmpc_distribute_static_init_4u(IdentTy *loc, int32_t global_tid,
uint32_t *plower, uint32_t *pupper,
int32_t *pstride, int32_t incr,
int32_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
@@ -634,7 +616,6 @@ void __kmpc_distribute_static_init_8(IdentTy *loc, int32_t global_tid,
int64_t *plower, int64_t *pupper,
int64_t *pstride, int64_t incr,
int64_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
@@ -645,19 +626,14 @@ void __kmpc_distribute_static_init_8u(IdentTy *loc, int32_t global_tid,
uint64_t *plower, uint64_t *pupper,
int64_t *pstride, int64_t incr,
int64_t chunk) {
- FunctionTracingRAII();
omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
mapping::isSPMDMode());
}
-void __kmpc_for_static_fini(IdentTy *loc, int32_t global_tid) {
- FunctionTracingRAII();
-}
+void __kmpc_for_static_fini(IdentTy *loc, int32_t global_tid) {}
-void __kmpc_distribute_static_fini(IdentTy *loc, int32_t global_tid) {
- FunctionTracingRAII();
-}
+void __kmpc_distribute_static_fini(IdentTy *loc, int32_t global_tid) {}
}
#pragma omp end declare target
More information about the Openmp-commits
mailing list