Fixed passing unpinned memory to async memcpy
This commit is contained in:
+27
-10
@@ -68,7 +68,7 @@ THE SOFTWARE.
|
|||||||
static const int release = 1;
|
static const int release = 1;
|
||||||
|
|
||||||
|
|
||||||
int HIP_LAUNCH_BLOCKING = 0;
|
int HIP_LAUNCH_BLOCKING = 1;
|
||||||
|
|
||||||
int HIP_PRINT_ENV = 0;
|
int HIP_PRINT_ENV = 0;
|
||||||
int HIP_TRACE_API= 0;
|
int HIP_TRACE_API= 0;
|
||||||
@@ -82,8 +82,8 @@ std::vector<int> g_hip_visible_devices; /* vector of integers that contains the
|
|||||||
|
|
||||||
//---
|
//---
|
||||||
// Chicken bits for disabling functionality to work around potential issues:
|
// Chicken bits for disabling functionality to work around potential issues:
|
||||||
int HIP_DISABLE_HW_KERNEL_DEP = 1;
|
int HIP_DISABLE_HW_KERNEL_DEP = 0;
|
||||||
int HIP_DISABLE_HW_COPY_DEP = 1;
|
int HIP_DISABLE_HW_COPY_DEP = 0;
|
||||||
|
|
||||||
int HIP_DISABLE_BIDIR_MEMCPY = 0;
|
int HIP_DISABLE_BIDIR_MEMCPY = 0;
|
||||||
int HIP_ONESHOT_COPY_DEP = 1; // TODO - setting this =1 is a good thing, reduces input deps
|
int HIP_ONESHOT_COPY_DEP = 1; // TODO - setting this =1 is a good thing, reduces input deps
|
||||||
@@ -2483,6 +2483,8 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
|
|||||||
|
|
||||||
stream = ihipSyncAndResolveStream(stream);
|
stream = ihipSyncAndResolveStream(stream);
|
||||||
|
|
||||||
|
bool trueAsync = true;
|
||||||
|
|
||||||
#if USE_AM_TRACKER
|
#if USE_AM_TRACKER
|
||||||
if (stream) {
|
if (stream) {
|
||||||
ihipDevice_t *device = stream->getDevice();
|
ihipDevice_t *device = stream->getDevice();
|
||||||
@@ -2501,13 +2503,17 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
|
|||||||
memcpy(dst, src, sizeBytes);
|
memcpy(dst, src, sizeBytes);
|
||||||
|
|
||||||
} else {
|
} else {
|
||||||
if (kind == hipMemcpyDefault) {
|
hc::accelerator acc;
|
||||||
hc::accelerator acc;
|
hc::AmPointerInfo dstAm(NULL, NULL, 0, acc, 0, 0);
|
||||||
hc::AmPointerInfo dstAm(NULL, NULL, 0, acc, 0, 0);
|
hc::AmPointerInfo srcAm(NULL, NULL, 0, acc, 0, 0);
|
||||||
hc::AmPointerInfo srcAm(NULL, NULL, 0, acc, 0, 0);
|
am_status_t statDst = hc::am_memtracker_getinfo(&dstAm, dst);
|
||||||
am_status_t statDst = hc::am_memtracker_getinfo(&dstAm, dst);
|
am_status_t statSrc = hc::am_memtracker_getinfo(&srcAm, src);
|
||||||
am_status_t statSrc = hc::am_memtracker_getinfo(&srcAm, src);
|
|
||||||
|
|
||||||
|
if(dstAm._appAllocationFlags != 1 || srcAm._appAllocationFlags != 1){
|
||||||
|
trueAsync = false;
|
||||||
|
}
|
||||||
|
|
||||||
|
if (kind == hipMemcpyDefault) {
|
||||||
if(statDst == AM_SUCCESS && statSrc == AM_SUCCESS){
|
if(statDst == AM_SUCCESS && statSrc == AM_SUCCESS){
|
||||||
if(dstAm._devicePointer != NULL){
|
if(dstAm._devicePointer != NULL){
|
||||||
if(srcAm._devicePointer != NULL){
|
if(srcAm._devicePointer != NULL){
|
||||||
@@ -2534,7 +2540,10 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
|
|||||||
hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1);
|
hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1);
|
||||||
|
|
||||||
ihipCommand_t copyType;
|
ihipCommand_t copyType;
|
||||||
if ((kind == hipMemcpyHostToDevice) || (kind == hipMemcpyDeviceToDevice)) {
|
if (kind == hipMemcpyHostToDevice ){
|
||||||
|
copyType = ihipCommandCopyH2D;
|
||||||
|
|
||||||
|
}else if(kind == hipMemcpyDeviceToDevice) {
|
||||||
copyType = ihipCommandCopyH2D;
|
copyType = ihipCommandCopyH2D;
|
||||||
} else if (kind == hipMemcpyDeviceToHost) {
|
} else if (kind == hipMemcpyDeviceToHost) {
|
||||||
copyType = ihipCommandCopyD2H;
|
copyType = ihipCommandCopyD2H;
|
||||||
@@ -2543,7 +2552,11 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
|
|||||||
copyType = ihipCommandCopyD2H;
|
copyType = ihipCommandCopyD2H;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
#if USE_ROCR_V2
|
#if USE_ROCR_V2
|
||||||
|
if(trueAsync == true){
|
||||||
|
|
||||||
hsa_signal_t depSignal;
|
hsa_signal_t depSignal;
|
||||||
int depSignalCnt = stream->copyCommand(ihip_signal, &depSignal, copyType);
|
int depSignalCnt = stream->copyCommand(ihip_signal, &depSignal, copyType);
|
||||||
|
|
||||||
@@ -2566,6 +2579,10 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
|
|||||||
// TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
|
// TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
|
||||||
e = hipErrorInvalidValue;
|
e = hipErrorInvalidValue;
|
||||||
}
|
}
|
||||||
|
}
|
||||||
|
else {
|
||||||
|
ihipSyncCopy(stream, dst, src, sizeBytes, kind);
|
||||||
|
}
|
||||||
}
|
}
|
||||||
} else {
|
} else {
|
||||||
e = hipErrorInvalidValue;
|
e = hipErrorInvalidValue;
|
||||||
|
|||||||
Reference in New Issue
Block a user