AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-283515 - Fix crashing in kernel launch on MGPUS

Просмотр исходного кода 
Fix wrong mixing of current device and stream device in
ihipModuleLaunchKernel() and hipLaunchCooperativeKernel().

Fix missing hipSetDevice() in hipMemcpyWithStream* tests.

Change-Id: I09333bb40d239bb42c832df5ea16d17eeaeff5e7
Этот коммит содержится в:
Tao Sang
2021-04-27 19:03:22 -04:00
коммит произвёл Tao Sang
родитель 6b14050424
Коммит a95ff95bf7
7 изменённых файлов: 25 добавлений и 20 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
rocclr/hip_internal.hpp
+3
Просмотреть файл
@@ -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<hip::Device*> 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);
rocclr/hip_module.cpp
+7 -5
Просмотреть файл
@@ -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<size_t>(gridDim.x) * blockDim.x;
size_t globalWorkSizeY = static_cast<size_t>(gridDim.y) * blockDim.y;
rocclr/hip_platform.cpp
+1 -7
Просмотреть файл
@@ -621,14 +621,8 @@ hipError_t ihipLaunchKernel(const void* hostFunction,
hipEvent_t stopEvent,
int flags)
{
hip::Stream* s = reinterpret_cast<hip::Stream*>(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);
rocclr/hip_stream.cpp
+7
Просмотреть файл
@@ -114,6 +114,13 @@ int Stream::DeviceId() const {
return device_->deviceId();
}
int Stream::DeviceId(const hipStream_t hStream) {
hip::Stream* s = reinterpret_cast<hip::Stream*>(hStream);
int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice();
assert(deviceId >= 0 && deviceId < static_cast<int>(g_devices.size()));
return deviceId;
}
void Stream::syncNonBlockingStreams() {
amd::ScopedLock lock(streamSetLock);
for (auto& it : streamSet) {
tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp
+4 -4
Просмотреть файл
@@ -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<const int*>(A_d[i]), static_cast<const int*>(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);
tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp
+2 -4
Просмотреть файл
@@ -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<const int*>(A_d[i]),
static_cast<const int*>(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.
tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp
+1
Просмотреть файл
@@ -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");