AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

SWDEV-294248 (Fixing Race Conditions):

Bläddra i källkod 
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


[ROCm/clr commit: 0670e49806]
This commit is contained in:
Ammar ELWazir
2021-08-07 01:04:45 -04:00
committad av Ammar Elwazir
förälder 0bda91b6ba
incheckning efc02f62d0
1 ändrade filer med 14 tillägg och 23 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_prof_api.h
+14 -23
Visa fil
@@ -188,42 +188,32 @@ template <int cid_>
class api_callbacks_spawner_t {
public:
api_callbacks_spawner_t() :
api_data_(NULL), was_enabled_on_construction_(is_enabled())
api_data_(NULL)
{
if (!was_enabled_on_construction_) return;
if (!is_enabled()) return;
if (cid_ < HIP_API_ID_FIRST || cid_ > HIP_API_ID_LAST) {
fprintf(stderr, "HIP %s bad id %d\n", __FUNCTION__, cid_);
abort();
}
static_assert(cid_ >= HIP_API_ID_FIRST || cid_ <= HIP_API_ID_LAST);
callbacks_table.sem_sync(cid_);
auto &entry = this->entry(cid_);
fun_ = std::make_pair(entry.fun, entry.arg);
act_ = std::make_pair(entry.act, entry.a_arg);
callbacks_table.sem_release(cid_);
hip_act_callback_t act = entry(cid_).act;
if (act != NULL) api_data_ = (hip_api_data_t*) act(cid_, NULL, NULL, NULL);
if (act_.first != NULL) api_data_ = (hip_api_data_t*) act_.first(cid_, NULL, NULL, NULL);
}
void call() {
hip_api_callback_t fun = entry(cid_).fun;
void* arg = entry(cid_).arg;
if (fun != NULL) {
fun(HIP_DOMAIN_ID, cid_, api_data_, arg);
if (fun_.first != NULL) {
fun_.first(HIP_DOMAIN_ID, cid_, api_data_, fun_.second);
api_data_->phase = ACTIVITY_API_PHASE_EXIT;
}
}
~api_callbacks_spawner_t() {
if (!was_enabled_on_construction_) return;
if (api_data_ != NULL) {
hip_api_callback_t fun = entry(cid_).fun;
void* arg = entry(cid_).arg;
hip_act_callback_t act = entry(cid_).act;
void* a_arg = entry(cid_).a_arg;
if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, api_data_, arg);
if (act != NULL) act(cid_, NULL, NULL, a_arg);
if (fun_.first != NULL) fun_.first(HIP_DOMAIN_ID, cid_, api_data_, fun_.second);
if (act_.first != NULL) act_.first(cid_, NULL, NULL, act_.second);
}
callbacks_table.sem_release(cid_);
}
hip_api_data_t* get_api_data_ptr() {
@@ -239,8 +229,9 @@ class api_callbacks_spawner_t {
return callbacks_table.entry(id);
}
std::pair<hip_api_callback_t, void *> fun_;
std::pair<hip_act_callback_t, void *> act_;
hip_api_data_t* api_data_;
bool was_enabled_on_construction_ = false;
};
template <>