Merge 'master' into 'amd-master'
Change-Id: Ib1251a94fa5c1d242d2cdd89341fbdc33644cf41
[ROCm/clr commit: 4783e4c2f3]
此提交包含在：
Jenkins
@@ -0,0 +1,20 @@
|
||||
---
|
||||
Language: Cpp
|
||||
BasedOnStyle: Google
|
||||
AlignEscapedNewlinesLeft: false
|
||||
ColumnLimit: 100
|
||||
DerivePointerAlignment: false
|
||||
IndentWrappedFunctionNames: false
|
||||
MaxEmptyLinesToKeep: 2
|
||||
SortIncludes: false
|
||||
IndentWidth: 4
|
||||
---
|
||||
Language: ObjC
|
||||
BasedOnStyle: Google
|
||||
AlignEscapedNewlinesLeft: false
|
||||
ColumnLimit: 100
|
||||
DerivePointerAlignment: false
|
||||
IndentWrappedFunctionNames: false
|
||||
MaxEmptyLinesToKeep: 2
|
||||
SortIncludes: false
|
||||
IndentWidth: 4
|
||||
@@ -220,8 +220,11 @@ endif()
|
||||
file(WRITE "${PROJECT_BINARY_DIR}/.hipVersion" ${_versionInfo})
|
||||
|
||||
# Build doxygen documentation
|
||||
add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMAKE_CURRENT_SOURCE_DIR}/docs/doxygen-input/doxy.cfg
|
||||
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/docs)
|
||||
find_program(DOXYGEN_EXE doxygen)
|
||||
if(DOXYGEN_EXE)
|
||||
add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} ${DOXYGEN_EXE} ${CMAKE_CURRENT_SOURCE_DIR}/docs/doxygen-input/doxy.cfg
|
||||
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/docs)
|
||||
endif()
|
||||
|
||||
#############################
|
||||
# Install steps
|
||||
@@ -364,8 +367,23 @@ endif()
|
||||
#############################
|
||||
# Code analysis
|
||||
#############################
|
||||
# Target: static_check
|
||||
add_custom_target(static_check COMMAND cppcheck --force --quiet --enable=warning,performance,portability,information,missingInclude src include -I /opt/rocm/include/hcc -I /opt/rocm/include --suppress=*:/opt/rocm/include/hcc/hc.hpp WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
|
||||
# Target: cppcheck
|
||||
find_program(CPPCHECK_EXE cppcheck)
|
||||
if(CPPCHECK_EXE)
|
||||
add_custom_target(cppcheck COMMAND ${CPPCHECK_EXE} --force --quiet --enable=warning,performance,portability,information,missingInclude src include -I /opt/rocm/include/hcc -I /opt/rocm/include --suppress=*:/opt/rocm/include/hcc/hc.hpp
|
||||
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
|
||||
endif()
|
||||
|
||||
#############################
|
||||
# Code formatting
|
||||
#############################
|
||||
# Target: clangformat
|
||||
find_program(CLANGFORMAT_EXE clang-format PATHS ${HCC_HOME}/bin)
|
||||
if(CLANGFORMAT_EXE)
|
||||
file(GLOB_RECURSE FORMAT_SOURCE_FILE_LIST *.cpp *.hpp *.h)
|
||||
add_custom_target(clangformat COMMAND ${CLANGFORMAT_EXE} -style=file -i ${FORMAT_SOURCE_FILE_LIST}
|
||||
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
|
||||
endif()
|
||||
|
||||
#############################
|
||||
# Testing steps
|
||||
|
||||
@@ -54,7 +54,7 @@ A stronger system-level fence can be specified when the event is created with hi
|
||||
- HIP/ROCm also supports the ability to cache host memory in the GPU using the "Non-Coherent" host memory allocations. This can provide performance benefit, but care must be taken to use the correct synchronization.
|
||||
|
||||
|
||||
## Unpinned Memory Transfer Optimizations
|
||||
## Unpinned Memory Transfer Optimization
|
||||
Please note that this document lists possible ways for experimenting with HIP stack to gain performance. Performance may vary from platform to platform.
|
||||
|
||||
### On Small BAR Setup
|
||||
@@ -79,11 +79,20 @@ stage the copy through an optimized pinned staging buffer, to implement H2D and
|
||||
|
||||
PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA engine.
|
||||
|
||||
By default staging buffers are used for unpinned memory transfers. Environment variables allow control over the unpinned copy algorithm and parameters:
|
||||
Unpinned memory transfer mode can be controlled using environment variable HCC_UNPINNED_COPY_MODE.
|
||||
|
||||
- HIP_PININPLACE - This environment variable forces the use of PinInPlace logic for all unpinned memory copies
|
||||
By default HCC_UNPINNED_COPY_MODE is set to 0, which uses default threshold values to decide which transfer way to use based on data size.
|
||||
|
||||
- HIP_OPTIMAL_MEM_TRANSFER- This environment variable enables a hybrid memory copy logic based on thresholds. These thresholds can be managed with following environment variables:
|
||||
- HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE - Threshold in bytes for H2D copy. For sizes smaller than threshold staging buffers logic would be used else PinInPlace logic.
|
||||
- HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING - Threshold in bytes for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic.
|
||||
- HIP_D2H_MEM_TRANSFER_THRESHOLD - Threshold in bytes for D2H copy. For sizes smaller than threshold staging buffer logic would be used else PinInPlace logic.
|
||||
Setting HCC_UNPINNED_COPY_MODE = 1, forces all unpinned transfer to use PinInPlace logic.
|
||||
|
||||
Setting HCC_UNPINNED_COPY_MODE = 2, forces all unpinned transfer to use Staging buffers.
|
||||
|
||||
Setting HCC_UNPINNED_COPY_MODE = 3, forces all unpinned transfer to use direct memcpy on large BAR systems.
|
||||
|
||||
Following environment variables can be used to control the transfer thresholds:
|
||||
|
||||
- HCC_H2D_STAGING_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic. By default it is set to 64.
|
||||
|
||||
- HCC_H2D_PININPLACE_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold staging buffers logic would be used else PinInPlace logic. By default it is set to 4096.
|
||||
|
||||
- HCC_D2H_PININPLACE_THRESHOLD - Threshold in KB for D2H copy. For sizes smaller than threshold staging buffer logic would be used else PinInPlace logic. By default it is set to 1024.
|
||||
|
||||
@@ -1420,7 +1420,56 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
|
||||
hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset,
|
||||
const void* src, size_t count, hipMemcpyKind kind);
|
||||
|
||||
/**
|
||||
* @brief Copies data between host and device.
|
||||
*
|
||||
* @param[in] dst Destination memory address
|
||||
* @param[in] srcArray Source memory address
|
||||
* @param[in] woffset Source starting X offset
|
||||
* @param[in] hOffset Source starting Y offset
|
||||
* @param[in] count Size in bytes to copy
|
||||
* @param[in] kind Type of transfer
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
|
||||
*/
|
||||
hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
|
||||
size_t count, hipMemcpyKind kind);
|
||||
|
||||
/**
|
||||
* @brief Copies data between host and device.
|
||||
*
|
||||
* @param[in] dst Destination memory address
|
||||
* @param[in] srcArray Source array
|
||||
* @param[in] srcoffset Offset in bytes of source array
|
||||
* @param[in] count Size of memory copy in bytes
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
|
||||
*/
|
||||
hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count);
|
||||
|
||||
/**
|
||||
* @brief Copies data between host and device.
|
||||
*
|
||||
* @param[in] dstArray Destination memory address
|
||||
* @param[in] dstOffset Offset in bytes of destination array
|
||||
* @param[in] srcHost Source host pointer
|
||||
* @param[in] count Size of memory copy in bytes
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
|
||||
*/
|
||||
hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count);
|
||||
|
||||
/**
|
||||
* @brief Copies data between host and device.
|
||||
*
|
||||
* @param[in] p 3D memory copy parameters
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
|
||||
*/
|
||||
hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p);
|
||||
|
||||
// doxygen end Memory
|
||||
|
||||
@@ -69,18 +69,16 @@ namespace hip_impl
|
||||
}
|
||||
};
|
||||
|
||||
using RAII_global = std::unique_ptr<void, decltype(hsa_amd_memory_unlock)*>;
|
||||
|
||||
const std::unordered_map<
|
||||
hsa_agent_t, std::vector<hsa_executable_t>>& executables();
|
||||
const std::unordered_map<
|
||||
std::uintptr_t,
|
||||
std::vector<std::pair<hsa_agent_t, Kernel_descriptor>>>& functions();
|
||||
const std::unordered_map<std::uintptr_t, std::string>& function_names();
|
||||
std::unordered_map<std::string, RAII_global>& globals();
|
||||
std::unordered_map<std::string, void*>& globals();
|
||||
|
||||
hsa_executable_t load_executable(
|
||||
const std::string& file,
|
||||
hsa_executable_t executable,
|
||||
hsa_agent_t agent);
|
||||
} // Namespace hip_impl.
|
||||
} // Namespace hip_impl.
|
||||
|
||||
@@ -483,6 +483,18 @@ inline static hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t
|
||||
return hipCUDAErrorTohipError(cudaMemcpyToArray(dst, wOffset, hOffset, src, count, hipMemcpyKindToCudaMemcpyKind(kind)));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind) {
|
||||
return hipCUDAErrorTohipError(cudaMemcpyFromArray(dst, srcArray, wOffset, hOffset, count, hipMemcpyKindToCudaMemcpyKind(kind)));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) {
|
||||
return hipCUResultTohipError(cuMemcpyAtoH(dst, (CUarray)srcArray, srcOffset, count));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) {
|
||||
return hipCUResultTohipError(cuMemcpyHtoA((CUarray)dstArray, dstOffset, srcHost, count));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceSynchronize() {
|
||||
return hipCUDAErrorTohipError(cudaDeviceSynchronize());
|
||||
}
|
||||
|
||||
@@ -453,6 +453,7 @@ hipError_t hipArrayCreate ( hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAlloc
|
||||
array[0]->width = pAllocateArray->width;
|
||||
array[0]->height = pAllocateArray->height;
|
||||
array[0]->isDrv = true;
|
||||
array[0]->textureType = hipTextureType2D;
|
||||
void ** ptr = &array[0]->data;
|
||||
if (ctx) {
|
||||
const unsigned am_flags = 0;
|
||||
@@ -1411,6 +1412,65 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset,
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
|
||||
size_t count, hipMemcpyKind kind) {
|
||||
|
||||
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind);
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((char *)dst, (char*)srcArray->data + wOffset, count, kind);
|
||||
}
|
||||
catch (ihipException &ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count)
|
||||
{
|
||||
HIP_INIT_SPECIAL_API((TRACE_MCMD), dstArray, dstOffset, srcHost, count);
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
try {
|
||||
stream->locked_copySync((char *)dstArray->data + dstOffset, srcHost, count, hipMemcpyHostToDevice);
|
||||
} catch (ihipException &ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count)
|
||||
{
|
||||
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, srcOffset, count);
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((char *)dst, (char*)srcArray->data + srcOffset, count, hipMemcpyDeviceToHost);
|
||||
}
|
||||
catch (ihipException &ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)
|
||||
{
|
||||
HIP_INIT_SPECIAL_API((TRACE_MCMD), p);
|
||||
|
||||
@@ -568,7 +568,6 @@ hipError_t hipModuleGetTexRef(
|
||||
const auto it = globals().find(name);
|
||||
if (it == globals().end()) return ihipLogStatus(hipErrorInvalidValue);
|
||||
|
||||
*texRef = static_cast<textureReference*>(it->second.get());
|
||||
|
||||
*texRef = reinterpret_cast<textureReference*>(it->second);
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
@@ -169,7 +169,7 @@ namespace
|
||||
lock_guard<mutex> lck{mtx};
|
||||
|
||||
if (globals().find(x) != globals().cend()) return;
|
||||
|
||||
globals().emplace(x, (void*)(it1->second.first));
|
||||
void* p = nullptr;
|
||||
hsa_amd_memory_lock(
|
||||
reinterpret_cast<void*>(it1->second.first),
|
||||
@@ -181,7 +181,6 @@ namespace
|
||||
hsa_executable_agent_global_variable_define(
|
||||
executable, agent, x.c_str(), p);
|
||||
|
||||
globals().emplace(x, RAII_global{p, hsa_amd_memory_unlock});
|
||||
}
|
||||
}
|
||||
|
||||
@@ -462,9 +461,9 @@ namespace hip_impl
|
||||
return r;
|
||||
}
|
||||
|
||||
unordered_map<string, RAII_global>& globals()
|
||||
unordered_map<string, void*>& globals()
|
||||
{
|
||||
static unordered_map<string, RAII_global> r;
|
||||
static unordered_map<string, void*> r;
|
||||
static once_flag f;
|
||||
call_once(f, []() { r.reserve(symbol_addresses().size()); });
|
||||
|
||||
@@ -491,4 +490,16 @@ namespace hip_impl
|
||||
|
||||
return executable;
|
||||
}
|
||||
} // Namespace hip_impl.
|
||||
|
||||
// To force HIP to load the kernels and to setup the function
|
||||
// symbol map on program startup
|
||||
class startup_kernel_loader {
|
||||
private:
|
||||
startup_kernel_loader() { functions(); }
|
||||
startup_kernel_loader(const startup_kernel_loader&) = delete;
|
||||
startup_kernel_loader& operator= (const startup_kernel_loader&) = delete;
|
||||
static startup_kernel_loader skl;
|
||||
};
|
||||
startup_kernel_loader startup_kernel_loader::skl;
|
||||
|
||||
} // Namespace hip_impl.
|
||||
|
||||
@@ -24,9 +24,9 @@ THE SOFTWARE.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* RUN: %t EXCLUDE_HIP_PLATFORM all
|
||||
* RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM all
|
||||
* RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM all
|
||||
* RUN: %t EXCLUDE_HIP_PLATFORM hcc
|
||||
* RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc
|
||||
* RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
新增問題並參考
封鎖使用者