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));