diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index f5eac616e4..f624d93eca 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -158,6 +158,8 @@ namespace hip { void Finish() const; /// Get device ID associated with the current stream; int DeviceId() const; + /// Get device ID associated with a stream; + static int DeviceId(const hipStream_t hStream); /// Returns if stream is null stream bool Null() const { return null_; } /// Returns the lock object for the current stream @@ -263,6 +265,7 @@ extern void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null extern std::vector g_devices; extern hipError_t ihipDeviceGetCount(int* count); extern int ihipGetDevice(); + extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags); extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); extern amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size); diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 6b61b2c6ff..6d5f902cb9 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -213,7 +213,7 @@ inline hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWork uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, uint32_t sharedMemBytes, - void** kernelParams, void** extra, + void** kernelParams, void** extra, int deviceId, uint32_t params = 0) { if (f == nullptr) { LogPrintfError("%s", "Function passed is null"); @@ -235,7 +235,8 @@ inline hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWork return hipErrorNotInitialized; } } - const amd::Device* device = hip::getCurrentDevice()->devices()[0]; + + const amd::Device* device = g_devices[deviceId]->devices()[0]; // Make sure dispatch doesn't exceed max workgroup size limit if (blockDimX * blockDimY * blockDimZ > device->info().maxWorkGroupSize_) { return hipErrorInvalidConfiguration; @@ -354,7 +355,8 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params); - HIP_RETURN_ONFAIL(PlatformState::instance().initStatManagedVarDevicePtr(ihipGetDevice())); + int deviceId = hip::Stream::DeviceId(hStream); + HIP_RETURN_ONFAIL(PlatformState::instance().initStatManagedVarDevicePtr(deviceId)); if (f == nullptr) { LogPrintfError("%s", "Function passed is null"); return hipErrorInvalidImage; @@ -365,7 +367,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipError_t status = ihipLaunchKernel_validate(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, blockDimX, - blockDimY, blockDimZ, sharedMemBytes, kernelParams, extra, params); + blockDimY, blockDimZ, sharedMemBytes, kernelParams, extra, deviceId, params); if (status != hipSuccess) { return status; } @@ -505,8 +507,8 @@ hipError_t hipLaunchCooperativeKernel(const void* f, HIP_INIT_API(hipLaunchCooperativeKernel, f, gridDim, blockDim, sharedMemBytes, hStream); - int deviceId = ihipGetDevice(); hipFunction_t func = nullptr; + int deviceId = hip::Stream::DeviceId(hStream); HIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, f, deviceId)); size_t globalWorkSizeX = static_cast(gridDim.x) * blockDim.x; size_t globalWorkSizeY = static_cast(gridDim.y) * blockDim.y; diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 2df8f5f028..01390fd9e3 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -621,14 +621,8 @@ hipError_t ihipLaunchKernel(const void* hostFunction, hipEvent_t stopEvent, int flags) { - hip::Stream* s = reinterpret_cast(stream); - int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice(); - if (deviceId == -1) { - LogPrintfError("Wrong Device Id: %d \n", deviceId); - HIP_RETURN(hipErrorNoDevice); - } - hipFunction_t func = nullptr; + int deviceId = hip::Stream::DeviceId(stream); hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId); if ((hip_error != hipSuccess) || (func == nullptr)) { HIP_RETURN(hipErrorInvalidDeviceFunction); diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp index b80fa28e17..50ef6960a9 100755 --- a/rocclr/hip_stream.cpp +++ b/rocclr/hip_stream.cpp @@ -114,6 +114,13 @@ int Stream::DeviceId() const { return device_->deviceId(); } +int Stream::DeviceId(const hipStream_t hStream) { + hip::Stream* s = reinterpret_cast(hStream); + int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice(); + assert(deviceId >= 0 && deviceId < static_cast(g_devices.size())); + return deviceId; +} + void Stream::syncNonBlockingStreams() { amd::ScopedLock lock(streamSetLock); for (auto& it : streamSet) { diff --git a/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp b/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp index 1dcf9a3d16..da46d51961 100755 --- a/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp +++ b/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp @@ -394,9 +394,11 @@ void HipMemcpyWithStreamtests::TestkindDefaultForDtoD(void) { int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices]; // Initialize and create the host and device elements for first device + HIPCHECK(hipSetDevice(0)); HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false); for (int i=1; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); HIPCHECK(hipMalloc(&A_d[i], Nbytes)); HIPCHECK(hipMalloc(&B_d[i], Nbytes)); HIPCHECK(hipMalloc(&C_d[i], Nbytes)); @@ -406,27 +408,25 @@ void HipMemcpyWithStreamtests::TestkindDefaultForDtoD(void) { hipStream_t stream[numDevices]; for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); HIPCHECK(hipStreamCreate(&stream[i])); } - HIPCHECK(hipSetDevice(0)); HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0])); HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0])); for (int i=1; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i])); HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDefault, stream[i])); } for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i], static_cast(A_d[i]), static_cast(B_d[i]), C_d[i], N); } for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipSetDevice(i)); // hipMemcpy will be on this device HIPCHECK(hipStreamSynchronize(stream[i])); HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N); diff --git a/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp b/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp index eecc7dd467..b5bb0a986e 100644 --- a/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp +++ b/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp @@ -463,6 +463,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) { int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices]; // Initialize and create the host and device elements for first device + HIPCHECK(hipSetDevice(0)); HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false); @@ -481,7 +482,6 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) { HIPCHECK(hipStreamCreate(&stream[i])); } - HIPCHECK(hipSetDevice(0)); HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0])); HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, @@ -491,7 +491,6 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) { // using hipMemcpyDefault kind that is numDevices in the setup. // 1st GPU start numbering from 0,1,2..n etc. for (int i=1; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i])); HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, @@ -499,14 +498,13 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) { } for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i], static_cast(A_d[i]), static_cast(B_d[i]), C_d[i], N); } for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipSetDevice(i)); // hipMemcpy will be on this device HIPCHECK(hipStreamSynchronize(stream[i])); HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); // Output of each GPU is getting validated with input of 1st GPU. diff --git a/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp b/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp index 668088458d..6cec0298c7 100644 --- a/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp +++ b/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp @@ -126,6 +126,7 @@ int main(int argc, char *argv[]) printf ("info: checking result on bus 0x%2x %s\n", props.pciBusID, props.name); printf ("info: copy Device2Host\n"); + HIPCHECK(hipSetDevice(j)); HIPCHECK( hipMemcpy(C_h, C_d[j], Nbytes, hipMemcpyDeviceToHost)); printf ("info: check result\n");