Commit Graph

11 Commits

Author SHA1 Message Date
Laurent Morichetti 6f78083f2a Fix a race condition when setting the callbacks
When acquiring the reader's lock (sem_sync()...sem_release()), it is
possible for a writer to squeeze by if the reader goes into sync_wait().
The writer could re-acquire the entry.sync between entry.sem > 0 and
entry.sync = 0.

void sync_wait(uint32_t id) {
    sem_decrement(id);
    while (entry(id).sync.load()) {}
    // <--- HERE
    sem_increment(id);
}

This could result in both the reader and the writer accessing
{ callback, arg } at the same time, and the reader could read
inconsistent data, for example: { new callback, old arg }.

The solution is to re-test entry.sync when returning from sync_wait():

void sem_sync(uint32_t id) {
    sem_increment(id);
-   if (entry(id).sync.load() == true) sync_wait(id);
+   while (entry(id).sync.load() == true) sync_wait(id);
}

Change-Id: I22f74f4cb9a5f027aac8aa4ed3e633acc19df4b8
2022-08-01 20:03:10 -04:00
Laurent Morichetti 73ee1afb17 Fix a build failure on Windows
Terse static_assert is only available starting with c++17. Add the
message argument since HIP is not compiled using c++17.

Change-Id: I77bc7ddb635e50ac3e6744d73c8751dc8299087e
2021-08-11 21:25:49 -07:00
Ammar ELWazir 0670e49806 SWDEV-294248 (Fixing Race Conditions):
Fixing race conditions that happened when enabling trace-period feature on the following code:
#include <hip/hip_runtime.h>
__global__ void
kernel ()
{
}
int
main (int argc, char **argv)
{
  for (size_t i = 0; i < 10000; ++i) {
    hipLaunchKernelGGL (kernel, 1, 1, 0, 0);
    hipDeviceSynchronize ();
  }
  return 0;
}

Change-Id: I6d2407ad06ecd602b93f53970859e1f4b67072e8
2021-08-11 11:06:10 -04:00
agunashe d9d9e81acb SWDEV-293742 - Update copyrights end year for hipamd
Change-Id: I08f620f84563a9214b59f1b943ed091b67229eab
2021-07-09 12:08:39 -04:00
Christophe Paquot d208afcb36 SWDEV-281658 - Preserve the callback IDs enum ordering
This change addresses the rocprofiler and HIP backward compatibility
issues. Before this patch, each time the hip_prof_str.h header was
generated, the ordering of the callbacks IDs changed, causing
incompatibilities between tools compiled with the old header and
runtimes compiled with the new headers (or vice versa).

To make the API callback IDs stable, the previous version of the header
is read to extract the enum values so that the same values can be
assigned in the new header.

Also, to make diffing different versions of the hip_prof_str.h easier
to read, all other sections (types, macros, helper functions) are now
alphabetically ordered.

If an update to the checked-in hip_prof_str.h file is required, the
cmake build is aborted and a message printed on stderr. The build will
not be successful until the checked-in hip_prof_str.h and the generated
hip_prof_str.h match.

Change-Id: I38b920e601185f7365a76a6584df91a7e8a11798
2021-07-09 11:17:09 -04:00
Christophe Paquot 0f7a47a95f SWDEV-281658 - Fix compilation warnings
In file included from /extra/lmoriche/hip-vdi/hip/rocclr/hip_internal.hpp:25,
                 from /extra/lmoriche/hip-vdi/hip/rocclr/hip_hmm.cpp:22:
/extra/lmoriche/hip-vdi/hip/rocclr/hip_prof_api.h: In constructor ‘api_callbacks_table_t::api_callbacks_table_t()’:
/extra/lmoriche/hip-vdi/hip/rocclr/hip_prof_api.h:72:59: warning: ‘void* memset(void*, int, size_t)’ clearing an object of type ‘struct api_callbacks_table_t::hip_cb_table_t’ with no trivial copy-assignment; use value-initialization instead [-Wclass-memaccess]
   72 |      memset(&callbacks_table_, 0, sizeof(callbacks_table_));
      |                                                           ^
/extra/lmoriche/hip-vdi/hip/rocclr/hip_prof_api.h:67:10: note: ‘struct api_callbacks_table_t::hip_cb_table_t’ declared here
   67 |   struct hip_cb_table_t {
      |          ^~~~~~~~~~~~~~

Address the above warning by providing default initialization of the
api_callbacks_table_t::hip_cb_table_t class members.

Change-Id: I69ea7c390c28cf3f8aec57f23566d6a3061a0365
2021-06-18 07:03:35 -07:00
Aleksey 4f0bf535c4 SWDEV-284883 - QMCPACK hangs when using rocprof --trace-period
Workaraund fix race condition in api_callbacks_spawner_t destructor.

Change-Id: I77d61492dfa484c3e477805c6f314a00dbb6cc32
2021-06-17 00:24:56 -04:00
Rahul Garg 19c84bc604 ROCMOPS-1956 - Push restructured code to hipamd
hipamd will have AMD's ROCCLR based HIP backend implementation

Change-Id: Id7de9634519b4ce46fca71a1b61f3d5b1e3fc459
2021-06-07 21:42:44 +00:00
Tao Sang 1cba7ec965 Remove hip-hcc codes: Part one
Remove hip-hcc codes from hip code base
Simplify hip CMakeLists.txt to exclude hip-hcc
Simplify cmake cmd for hip-rocclr building
Some minor fixes

Change-Id: I1ae357ecfd638d6c25bca293c1724b026be21ecd
2020-12-09 15:49:47 -05:00
Evgeny ed29c10394 tracing layer unifying with hcc: removing API_ID_ANY 2019-10-08 19:09:43 -05:00
Evgeny c6600ba26b prof layer includes refactoring 2019-06-11 20:13:29 -05:00