SWDEV-273943 - Enable runtimeApi.event dtest

Add device_id_ in hip::event to match cuda behaviour in hipEventQuery() and hipEventRecord(). Enable hipEventElapsedTime test on AMD platform. Workarround sporadic crash of hipEventIpc test due to some bug of event ipc. Add missing hipEventDestroy() in some event tests. Fix some logic code errors. Fix typo in comment. Change-Id: I9ec74c475161b3e31df48d193449023e921f2924 [ROCm/clr commit: 629ea5fe19]
2021-03-09 12:00:16 -05:00
parent d09885f10d
commit e212e1a947
9 changed files with 69 additions and 14 deletions
@@ -23,7 +23,7 @@ ROCm defines two coherency options for host memory:
 IP provides the developer with controls to select which type of memory is used via allocation flags passed to hipHostMalloc and the HIP_HOST_COHERENT environment variable:
 - hipHostllocCoherent=0, hipHostMallocNonCoherent=0: Use HIP_HOST_COHERENT environment variable:
    - If HIP_HOST_COHERENT is 1 or undefined, the host memory allocation is coherent.
-    - If host memory is `defined and 0: the host memory allocation is non-coherent.
+    - If HIP_HOST_COHERENT is `defined and 0: the host memory allocation is non-coherent.
 - hipHostMallocCoherent=1, hipHostMallocNonCoherent=0: The host memory allocation will be coherent.  HIP_HOST_COHERENT env variable is ignored.
 - hipHostMallocCoherent=0, hipHostMallocNonCoherent=1: The host memory allocation will be non-coherent.  HIP_HOST_COHERENT env variable is ignored.
 - hipHostMallocCoherent=1, hipHostMallocNonCoherent=1: Illegal.
@@ -147,7 +147,7 @@ enum hipLimit_t {
                /// obtain more precise timings of commands between events.  The flag is a no-op on
                /// CUDA platforms.
 #define hipEventReleaseToSystem                                                                    \
-    0x80000000  /// < Use a system-scope release that when recording this event.  This flag is
+    0x80000000  /// < Use a system-scope release when recording this event.  This flag is
                /// useful to make non-coherent host memory visible to the host.  The flag is a
                /// no-op on CUDA platforms.

@@ -207,7 +207,6 @@ hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
    if (e == nullptr) {
      return hipErrorOutOfMemory;
    }
-
    *event = reinterpret_cast<hipEvent_t>(e);
  } else {
    return hipErrorInvalidValue;
@@ -258,9 +257,8 @@ hipError_t hipEventDestroy(hipEvent_t event) {
    if (!amd::Os::MemoryUnmapFile(e->ipc_evt_.ipc_shmem_,sizeof(hip::ihipIpcEventShmem_t))) {
      HIP_RETURN(hipErrorInvalidHandle);
    }
-  } else {
-    delete e;
  }
+  delete e;
  HIP_RETURN(hipSuccess);
 }

@@ -278,14 +276,23 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) {
  hip::Event* eStart = reinterpret_cast<hip::Event*>(start);
  hip::Event* eStop  = reinterpret_cast<hip::Event*>(stop);

+  if (eStart->deviceId() != eStop->deviceId()) {
+    HIP_RETURN(hipErrorInvalidHandle);
+  }
+
  HIP_RETURN(eStart->elapsedTime(*eStop, *ms), "Elapsed Time = ", *ms);
 }

 // ================================================================================================
 bool createIpcEventShmemIfNeeded(hip::Event::ihipIpcEvent_t& ipc_evt) {
 #if !defined(_MSC_VER)
+  if (ipc_evt.ipc_shmem_) {
+    // ipc_shmem_ already created, no need to create it again
+    return true;
+  }
  char name_template[] = "/tmp/eventXXXXXX";
  int temp_fd = mkstemp(name_template);
+
  ipc_evt.ipc_name_ = name_template;
  ipc_evt.ipc_name_.replace(0, 5, "/hip_");
  if (!amd::Os::MemoryMapFileTruncated(ipc_evt.ipc_name_.c_str(),
@@ -298,6 +305,8 @@ bool createIpcEventShmemIfNeeded(hip::Event::ihipIpcEvent_t& ipc_evt) {
  for (uint32_t sig_idx = 0; sig_idx < IPC_SIGNALS_PER_EVENT; ++sig_idx) {
    ipc_evt.ipc_shmem_->signal[sig_idx] = 0;
  }
+
+  close(temp_fd);
  return true;
 #else
  return false;
@@ -312,7 +321,13 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
  }

  hip::Event* e = reinterpret_cast<hip::Event*>(event);
+
  amd::HostQueue* queue = hip::getQueue(stream);
+
+  if (g_devices[e->deviceId()]->devices()[0] != &queue->device()) {
+    HIP_RETURN(hipErrorInvalidHandle);
+  }
+
  bool isRecorded = e->isRecorded();
  if ((e->flags & hipEventInterprocess) && !isRecorded) {
    amd::Command* command = queue->getLastQueuedCommand(true);
@@ -330,6 +345,8 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
    }
    // Lock signal.
    e->ipc_evt_.ipc_shmem_->signal[offset] = 1;
+    e->ipc_evt_.ipc_shmem_->owners_device_id = e->deviceId();
+
    std::atomic<int> *signal = &e->ipc_evt_.ipc_shmem_->signal[offset];
    StreamCallback* cbo = new StreamCallback(stream,
                    reinterpret_cast<hipStreamCallback_t> (ipcEventCallback), signal, command);
@@ -421,7 +438,10 @@ hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle)
                    (const void**) &(ipc_evt.ipc_shmem_), sizeof(hip::ihipIpcEventShmem_t))) {
    HIP_RETURN(hipErrorInvalidValue);
  }
+
  ipc_evt.ipc_shmem_->owners += 1;
+  e->setDeviceId(ipc_evt.ipc_shmem_->owners_device_id.load());
+
  HIP_RETURN(hipSuccess);
 #else
  assert(0 && "Unimplemented");
@@ -57,6 +57,7 @@ public:
 #define IPC_SIGNALS_PER_EVENT 32
 typedef struct ihipIpcEventShmem_s {
  std::atomic<int> owners;
+  std::atomic<int> owners_device_id;
  std::atomic<int> read_index;
  std::atomic<int> write_index;
  std::atomic<int> signal[IPC_SIGNALS_PER_EVENT];
@@ -67,6 +68,7 @@ public:
  Event(unsigned int flags) : flags(flags), lock_("hipEvent_t", true),
                              event_(nullptr), recorded_(false) {
    // No need to init event_ here as addMarker does that
+    device_id_ = hip::getCurrentDevice()->deviceId(); // Created in current device ctx
  }

  ~Event() {
@@ -84,8 +86,10 @@ public:
  void addMarker(amd::HostQueue* queue, amd::Command* command, bool record);
  bool isRecorded() { return recorded_; }
  amd::Monitor& lock() { return lock_; }
+  const int deviceId() { return device_id_; }
+  void setDeviceId(int id) { device_id_ = id; }

-   //IPC Events
+  //IPC Events
  struct ihipIpcEvent_t {
    std::string ipc_name_;
    int ipc_fd_;
@@ -101,7 +105,7 @@ private:
  amd::Monitor lock_;
  amd::HostQueue* stream_;
  amd::Event* event_;
-
+  int device_id_;
  //! Flag to indicate hipEventRecord has been called. This is needed except for
  //! hip*ModuleLaunchKernel API which takes start and stop events so no
  //! hipEventRecord is called. Cleanup needed once those APIs are deprecated.
@@ -68,6 +68,7 @@ bool hipEvent_Nullcheck() {
    printf("hipEventDestroy failed when nullptr is passed \n");
    TestStatus = false;
  }
+
  return TestStatus;
 }

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
 * TEST: %t
 * HIT_END
 */
@@ -49,6 +49,8 @@ void NegativeTests(){
        HIPCHECK(hipEventCreateWithFlags(&start,hipEventDisableTiming));
        HIPCHECK(hipEventCreateWithFlags(&stop,hipEventDisableTiming));
        HIPASSERT(hipEventElapsedTime(&timeElapsed, start, stop) == hipErrorInvalidHandle);
+        HIPCHECK(hipEventDestroy(start));
+        HIPCHECK(hipEventDestroy(stop));
    }

    // events created different devices
@@ -59,7 +61,9 @@ void NegativeTests(){
            // create event on dev=0
            HIPCHECK(hipSetDevice(0));
            hipEvent_t start;
+            hipEvent_t start1;
            HIPCHECK(hipEventCreate(&start));
+            HIPCHECK(hipEventCreate(&start1));

            HIPCHECK(hipEventRecord(start, nullptr));
            HIPCHECK(hipEventSynchronize(start));
@@ -69,11 +73,19 @@ void NegativeTests(){
            hipEvent_t stop;
            HIPCHECK(hipEventCreate(&stop));

+            // start1 on device 0 but null stream on device 1
+            HIPASSERT(hipEventRecord(start1, nullptr) == hipErrorInvalidHandle);
+
            HIPCHECK(hipEventRecord(stop, nullptr));
            HIPCHECK(hipEventSynchronize(stop));

            float tElapsed = 1.0f;
+            // start on device 0 but stop on device 1
            HIPASSERT(hipEventElapsedTime(&tElapsed,start,stop) == hipErrorInvalidHandle);
+
+            HIPCHECK(hipEventDestroy(start));
+            HIPCHECK(hipEventDestroy(start1));
+            HIPCHECK(hipEventDestroy(stop));
        }
    }
 }
@@ -92,7 +104,10 @@ void PositiveTest(){
    HIPCHECK(hipEventSynchronize(stop));

    float tElapsed = 1.0f;
-    HIPCHECK(hipEventElapsedTime(&tElapsed,start,stop));
+    HIPCHECK(hipEventElapsedTime(&tElapsed, start, stop));
+
+    HIPCHECK(hipEventDestroy(start));
+    HIPCHECK(hipEventDestroy(stop));
 }

 int main(){
@@ -78,7 +78,7 @@ int main(int argc, char* argv[]) {


        float eventMs = 1.0f;
-        // should fail
+        // should fail due to hipEventDisableTiming
        HIPASSERT(hipSuccess != hipEventElapsedTime(&eventMs, start, stop));
        float hostMs = HipTest::elapsed_time(hostStart, hostStop);

@@ -97,16 +97,27 @@ int main(int argc, char* argv[]) {
    HIPCHECK(hipEventSynchronize(ipc_event));

    HIPCHECK(hipEventDestroy(ipc_event));
+
+#ifndef __HIP_PLATFORM_AMD__
    HIPCHECK(hipEventDestroy(start));
    HIPCHECK(hipEventDestroy(stop));
+#endif

    HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));

-
    printf("check:\n");

+#ifdef __HIP_PLATFORM_AMD__
+    // Due to implementation bug of hipEventInterprocess, as a workaround,
+    // we have to move hipEventDestroy here. Otherwise sporadic crash will
+    // happen in above hipMemcpy(). If hipEventInterprocess is officially
+    // implemented, we should revisit here and move these back to match cuda
+    // event behavior.
+    HIPCHECK(hipEventDestroy(start));
+    HIPCHECK(hipEventDestroy(stop));
+#endif
+
    HipTest::checkVectorADD(A_h, B_h, C_h, N, true);

-
    passed();
 }
@@ -91,6 +91,8 @@ int main(int argc, char* argv[]) {

    HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));

+    HIPCHECK(hipEventDestroy(start));
+    HIPCHECK(hipEventDestroy(stop));

    printf("check:\n");

@@ -104,11 +104,10 @@ void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, hipStream_
            assert(0);
    };

-
    float t;

    hipError_t e = hipEventElapsedTime(&t, start, start);
-    if ((e != hipSuccess) && (e != hipErrorNotReady)) {
+    if ((e != hipSuccess) && (e != hipErrorNotReady || syncMode != syncNone)) {
        failed("start event not in expected state, was %d=%s\n", e, hipGetErrorName(e));
    }

@@ -143,6 +142,9 @@ void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, hipStream_
        HIPCHECK_API(hipEventElapsedTime(&t, start, neverRecorded), hipErrorInvalidHandle);
    }

+    HIPCHECK(hipEventDestroy(neverRecorded));
+    HIPCHECK(hipEventDestroy(timingDisabled));
+
    HIPCHECK(hipEventDestroy(start));
    HIPCHECK(hipEventDestroy(stop));