diff --git a/hipamd/src/hip_conversions.hpp b/hipamd/src/hip_conversions.hpp index 17866d10b4..11fe66d363 100644 --- a/hipamd/src/hip_conversions.hpp +++ b/hipamd/src/hip_conversions.hpp @@ -604,6 +604,7 @@ std::pair getMemoryType(const hipMemcpyKind kind) case hipMemcpyDeviceToHost: return {hipMemoryTypeDevice, hipMemoryTypeHost}; case hipMemcpyDeviceToDevice: + case hipMemcpyDeviceToDeviceNoCU: return {hipMemoryTypeDevice, hipMemoryTypeDevice}; case hipMemcpyDefault: return {hipMemoryTypeUnified, hipMemoryTypeUnified}; diff --git a/hipamd/src/hip_formatting.hpp b/hipamd/src/hip_formatting.hpp index c4e874bd02..3073d8dbd0 100644 --- a/hipamd/src/hip_formatting.hpp +++ b/hipamd/src/hip_formatting.hpp @@ -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"; }; diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index c47e82e8d2..eaef79bf5b 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -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; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 71f9db472d..e4fb54bd78 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -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; } diff --git a/hipamd/src/trace_helper.h b/hipamd/src/trace_helper.h index aed87dd214..9ae9c4ed4b 100644 --- a/hipamd/src/trace_helper.h +++ b/hipamd/src/trace_helper.h @@ -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); }; diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index 14b1827b47..daefffb2bd 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -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;