AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-232918

Browse Source 
hipEventRecord is much slower in hipclang/vdi

- Make sure default streams don't sync each other.
- Add null stream into the list of default streams.
- Code clean-up to simplify queue look-up.

Change-Id: I36e1fc8d86a600e3dce806694d95d146ed8afd03
This commit is contained in:
German Andryeyev
2020-04-23 16:54:48 -04:00
parent cf6df62b3a
commit 5a7c094729
6 changed files with 95 additions and 134 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/vdi/hip_platform.cpp
+2 -2
View File
@@ -733,7 +733,7 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
PlatformState::instance().popExec(exec);
hip::Stream* stream = reinterpret_cast<hip::Stream*>(exec.hStream_);
int deviceId = (stream != nullptr)? stream->device->deviceId() : ihipGetDevice();
int deviceId = (stream != nullptr)? stream->DeviceId() : ihipGetDevice();
if (deviceId == -1) {
DevLogPrintfError("Wrong DeviceId: %d \n", deviceId);
HIP_RETURN(hipErrorNoDevice);
@@ -1212,7 +1212,7 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
stream);
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
int deviceId = (s != nullptr)? s->device->deviceId() : ihipGetDevice();
int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice();
if (deviceId == -1) {
DevLogPrintfError("Wrong Device Id: %d \n", deviceId);
HIP_RETURN(hipErrorNoDevice);