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

SWDEV-408180 - Address possible cornercases

Browse Source 
- Address corner cases that can arise with the new
hipMemcpyDeviceToDeviceNoCU enum
- Better log

Change-Id: I6035b901f8d616741054b7a5ff4f67956329ac57
This commit is contained in:
Saleel Kudchadker
2023-10-20 18:00:02 +00:00
parent de2b06a7a7
commit 5662d4037c
6 changed files with 34 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_conversions.hpp
+1
View File
@@ -604,6 +604,7 @@ std::pair<hipMemoryType, hipMemoryType> getMemoryType(const hipMemcpyKind kind)
case hipMemcpyDeviceToHost:
return {hipMemoryTypeDevice, hipMemoryTypeHost};
case hipMemcpyDeviceToDevice:
case hipMemcpyDeviceToDeviceNoCU:
return {hipMemoryTypeDevice, hipMemoryTypeDevice};
case hipMemcpyDefault:
return {hipMemoryTypeUnified, hipMemoryTypeUnified};
hipamd/src/hip_formatting.hpp
+3
View File
@@ -86,6 +86,9 @@ inline std::ostream& operator<<(std::ostream& os, const hipMemcpyKind& s) {
case hipMemcpyDefault:
os << "hipMemcpyDefault";
break;
case hipMemcpyDeviceToDeviceNoCU:
os << "hipMemcpyDeviceToDeviceNoCU";
break;
default:
os << "hipMemcpyDefault";
};
hipamd/src/hip_graph_internal.hpp
+4 -2
View File
@@ -1535,7 +1535,8 @@ class GraphMemcpyNodeFromSymbol : public GraphMemcpyNode1D {
if (dstMemory == nullptr && kind != hipMemcpyDeviceToHost && kind != hipMemcpyDefault) {
return hipErrorInvalidMemcpyDirection;
} else if (dstMemory != nullptr && dstMemory->getMemFlags() == 0 &&
kind != hipMemcpyDeviceToDevice && kind != hipMemcpyDefault) {
kind != hipMemcpyDeviceToDevice && kind != hipMemcpyDeviceToDeviceNoCU
&& kind != hipMemcpyDefault) {
return hipErrorInvalidMemcpyDirection;
} else if (kind == hipMemcpyHostToHost || kind == hipMemcpyHostToDevice) {
return hipErrorInvalidMemcpyDirection;
@@ -1625,7 +1626,8 @@ class GraphMemcpyNodeToSymbol : public GraphMemcpyNode1D {
if (srcMemory == nullptr && kind != hipMemcpyHostToDevice && kind != hipMemcpyDefault) {
return hipErrorInvalidValue;
} else if (srcMemory != nullptr && srcMemory->getMemFlags() == 0 &&
kind != hipMemcpyDeviceToDevice && kind != hipMemcpyDefault) {
kind != hipMemcpyDeviceToDevice && kind != hipMemcpyDeviceToDeviceNoCU
&& kind != hipMemcpyDefault) {
return hipErrorInvalidValue;
} else if (kind == hipMemcpyHostToHost || kind == hipMemcpyDeviceToHost) {
return hipErrorInvalidValue;
hipamd/src/hip_memory.cpp
+22 -6
View File
@@ -337,6 +337,8 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
memObj->getUserData().deviceId = hip::getCurrentDevice()->deviceId();
return hipSuccess;
}
// ================================================================================================
bool IsHtoHMemcpyValid(void* dst, const void* src, hipMemcpyKind kind) {
size_t sOffset = 0;
amd::Memory* srcMemory = getMemoryObject(src, sOffset);
@@ -349,6 +351,8 @@ bool IsHtoHMemcpyValid(void* dst, const void* src, hipMemcpyKind kind) {
}
return true;
}
// ================================================================================================
hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind) {
if (dst == nullptr || src == nullptr) {
@@ -371,6 +375,7 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
return hipSuccess;
}
// ================================================================================================
hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hip::Stream& stream, bool isAsync) {
amd::Command::EventWaitList waitList;
@@ -465,6 +470,8 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
}
return hipSuccess;
}
// ================================================================================================
bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind) {
size_t sOffset = 0;
amd::Memory* srcMemory = getMemoryObject(src, sOffset);
@@ -477,10 +484,13 @@ bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind) {
}
return false;
}
// ================================================================================================
void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, hip::Stream& stream) {
stream.finish();
memcpy(dst, src, sizeBytes);
}
// ================================================================================================
hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hip::Stream& stream, bool isHostAsync, bool isGPUAsync) {
@@ -1281,8 +1291,10 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_RETURN(ihipMalloc(ptr, sizeBytes, flags), (ptr != nullptr)? *ptr : nullptr);
};
inline hipError_t ihipMemcpySymbol_validate(const void* symbol, size_t sizeBytes, size_t offset, size_t &sym_size, hipDeviceptr_t &device_ptr) {
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size));
inline hipError_t ihipMemcpySymbol_validate(const void* symbol, size_t sizeBytes,
size_t offset, size_t &sym_size, hipDeviceptr_t &device_ptr) {
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(),
&device_ptr, &sym_size));
/* Size Check to make sure offset is correct */
if ((offset + sizeBytes) > sym_size) {
@@ -1299,7 +1311,8 @@ hipError_t hipMemcpyToSymbol_common(const void* symbol, const void* src, size_t
size_t offset, hipMemcpyKind kind, hipStream_t stream=nullptr) {
CHECK_STREAM_CAPTURING();
if (kind != hipMemcpyHostToDevice && kind != hipMemcpyDeviceToDevice) {
if (kind != hipMemcpyHostToDevice && (kind != hipMemcpyDeviceToDevice ||
kind != hipMemcpyDeviceToDeviceNoCU)) {
HIP_RETURN(hipErrorInvalidMemcpyDirection);
}
@@ -1332,7 +1345,8 @@ hipError_t hipMemcpyFromSymbol_common(void* dst, const void* symbol, size_t size
size_t offset, hipMemcpyKind kind, hipStream_t stream=nullptr) {
CHECK_STREAM_CAPTURING();
if (kind != hipMemcpyDeviceToHost && kind != hipMemcpyDeviceToDevice) {
if (kind != hipMemcpyDeviceToHost && (kind != hipMemcpyDeviceToDevice ||
kind != hipMemcpyDeviceToDeviceNoCU)) {
HIP_RETURN(hipErrorInvalidMemcpyDirection);
}
@@ -1365,7 +1379,8 @@ hipError_t hipMemcpyToSymbolAsync_common(const void* symbol, const void* src, si
size_t offset, hipMemcpyKind kind, hipStream_t stream) {
STREAM_CAPTURE(hipMemcpyToSymbolAsync, stream, symbol, src, sizeBytes, offset, kind);
if (kind != hipMemcpyHostToDevice && kind != hipMemcpyDeviceToDevice) {
if (kind != hipMemcpyHostToDevice && (kind != hipMemcpyDeviceToDevice ||
kind != hipMemcpyDeviceToDeviceNoCU)) {
return hipErrorInvalidMemcpyDirection;
}
@@ -1397,7 +1412,8 @@ hipError_t hipMemcpyFromSymbolAsync_common(void* dst, const void* symbol, size_t
size_t offset, hipMemcpyKind kind, hipStream_t stream) {
STREAM_CAPTURE(hipMemcpyFromSymbolAsync, stream, dst, symbol, sizeBytes, offset, kind);
if (kind != hipMemcpyDeviceToHost && kind != hipMemcpyDeviceToDevice) {
if (kind != hipMemcpyDeviceToHost && (kind != hipMemcpyDeviceToDevice ||
kind != hipMemcpyDeviceToDeviceNoCU)) {
return hipErrorInvalidMemcpyDirection;
}
hipamd/src/trace_helper.h
+1
View File
@@ -199,6 +199,7 @@ inline std::string ToString(hipMemcpyKind v) {
CASE_STR(hipMemcpyDeviceToHost);
CASE_STR(hipMemcpyDeviceToDevice);
CASE_STR(hipMemcpyDefault);
CASE_STR(hipMemcpyDeviceToDeviceNoCU);
default:
return ToHexString(v);
};
rocclr/device/rocm/rocdevice.cpp
+3 -2
View File
@@ -1329,7 +1329,7 @@ bool Device::populateOCLDeviceConstants() {
}
}
}
gpuvm_segment_max_alloc_ =
uint64_t(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u);
assert(gpuvm_segment_max_alloc_ > 0);
@@ -2128,7 +2128,8 @@ void* Device::hostAlloc(size_t size, size_t alignment, MemorySegment mem_seg) co
assert(segment.handle != 0);
hsa_status_t stat = hsa_amd_memory_pool_allocate(segment, size, 0, &ptr);
ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Allocate hsa host memory %p, size 0x%zx", ptr, size);
ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Allocate hsa host memory %p, size 0x%zx,"
" numa_node = %d", ptr, size, preferred_numa_node_);
if (stat != HSA_STATUS_SUCCESS) {
LogPrintfError("Fail allocation host memory with err %d", stat);
return nullptr;