From e212e1a947c8076d7c855550acc8a3ef8db0c577 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Tue, 9 Mar 2021 12:00:16 -0500 Subject: [PATCH] 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: 629ea5fe194ff837ce54e7e40a6e4b1a25ff6f2a] --- .../docs/markdown/hip_programming_guide.md | 2 +- .../include/hip/amd_detail/hip_runtime_api.h | 2 +- projects/clr/hipamd/rocclr/hip_event.cpp | 26 ++++++++++++++++--- projects/clr/hipamd/rocclr/hip_event.hpp | 8 ++++-- .../tests/src/runtimeApi/event/hipEvent.cpp | 1 + .../runtimeApi/event/hipEventElapsedTime.cpp | 19 ++++++++++++-- .../src/runtimeApi/event/hipEventIpc.cpp | 17 +++++++++--- .../src/runtimeApi/event/hipEventRecord.cpp | 2 ++ .../src/runtimeApi/event/record_event.cpp | 6 +++-- 9 files changed, 69 insertions(+), 14 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_programming_guide.md b/projects/clr/hipamd/docs/markdown/hip_programming_guide.md index 6a85a06f47..617a5a407d 100644 --- a/projects/clr/hipamd/docs/markdown/hip_programming_guide.md +++ b/projects/clr/hipamd/docs/markdown/hip_programming_guide.md @@ -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. diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h index 20bab9e779..8ce3bdac5f 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h @@ -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. diff --git a/projects/clr/hipamd/rocclr/hip_event.cpp b/projects/clr/hipamd/rocclr/hip_event.cpp index fe42f9cef3..6633f5fe79 100755 --- a/projects/clr/hipamd/rocclr/hip_event.cpp +++ b/projects/clr/hipamd/rocclr/hip_event.cpp @@ -207,7 +207,6 @@ hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { if (e == nullptr) { return hipErrorOutOfMemory; } - *event = reinterpret_cast(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(start); hip::Event* eStop = reinterpret_cast(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(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 *signal = &e->ipc_evt_.ipc_shmem_->signal[offset]; StreamCallback* cbo = new StreamCallback(stream, reinterpret_cast (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"); diff --git a/projects/clr/hipamd/rocclr/hip_event.hpp b/projects/clr/hipamd/rocclr/hip_event.hpp index b26088fcc9..71a6c44b10 100644 --- a/projects/clr/hipamd/rocclr/hip_event.hpp +++ b/projects/clr/hipamd/rocclr/hip_event.hpp @@ -57,6 +57,7 @@ public: #define IPC_SIGNALS_PER_EVENT 32 typedef struct ihipIpcEventShmem_s { std::atomic owners; + std::atomic owners_device_id; std::atomic read_index; std::atomic write_index; std::atomic 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. diff --git a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEvent.cpp b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEvent.cpp index eacb7669cd..e9b3f4db07 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEvent.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEvent.cpp @@ -68,6 +68,7 @@ bool hipEvent_Nullcheck() { printf("hipEventDestroy failed when nullptr is passed \n"); TestStatus = false; } + return TestStatus; } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventElapsedTime.cpp b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventElapsedTime.cpp index d4207332e1..6010ce9ce3 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventElapsedTime.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventElapsedTime.cpp @@ -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(){ diff --git a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventIpc.cpp b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventIpc.cpp index f5aa92b719..bd8db4c9f2 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventIpc.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventIpc.cpp @@ -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(); } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventRecord.cpp b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventRecord.cpp index a2a9beaf79..8e80b46275 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventRecord.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/event/hipEventRecord.cpp @@ -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"); diff --git a/projects/clr/hipamd/tests/src/runtimeApi/event/record_event.cpp b/projects/clr/hipamd/tests/src/runtimeApi/event/record_event.cpp index d276d17d25..50102b5e70 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/event/record_event.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/event/record_event.cpp @@ -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));