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

P4 to Git Change 1981065 by cpaquot@cpaquot-ocl-lc-lnx on 2019/08/09 16:31:44

Просмотр исходного кода 
SWDEV-198556 - [HIP] Use src/dstMemory->getContext instead of host_context.
	Also relax the check for P2P copies in case of hipMemcpy(hostMalloced, hipMalloced(dev1), dev0)

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#67 edit


[ROCm/clr commit: 7bbbdabfb0]
Этот коммит содержится в:
foreman
2019-08-09 16:34:38 -04:00
родитель 5a0cdd92cb
Коммит e7c316bc92
1 изменённых файлов: 25 добавлений и 31 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/clr/hipamd/api/hip/hip_memory.cpp
+25 -31
Просмотреть файл
@@ -100,7 +100,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
} else if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
amd::HostQueue* pQueue = &queue;
if (queueDevice != dstMemory->getContext().devices()[0]) {
pQueue = hip::getNullStream(*hip::host_context);
pQueue = hip::getNullStream(dstMemory->getContext());
}
command = new amd::WriteMemoryCommand(*pQueue, CL_COMMAND_WRITE_BUFFER, waitList,
*dstMemory->asBuffer(), dOffset, sizeBytes, src);
@@ -108,43 +108,37 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
} else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
amd::HostQueue* pQueue = &queue;
if (queueDevice != srcMemory->getContext().devices()[0]) {
pQueue = hip::getNullStream(*hip::host_context);
pQueue = hip::getNullStream(srcMemory->getContext());
}
command = new amd::ReadMemoryCommand(*pQueue, CL_COMMAND_READ_BUFFER, waitList,
*srcMemory->asBuffer(), sOffset, sizeBytes, dst);
isAsync = false;
} else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
static const uint hostMem = CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR;
if ((kind == hipMemcpyDeviceToDevice ||
kind == hipMemcpyDefault) &&
((srcMemory->getMemFlags() & hostMem) == 0) &&
((dstMemory->getMemFlags() & hostMem) == 0)) {
if (queueDevice != srcMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
}
command->release();
return hipSuccess;
if (queueDevice != srcMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
}
if (queueDevice != dstMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
}
command->release();
return hipSuccess;
command->release();
return hipSuccess;
}
if (queueDevice != dstMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
}
command->release();
return hipSuccess;
}
command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);