diff --git a/projects/clr/hipamd/.clang-format b/projects/clr/hipamd/.clang-format new file mode 100644 index 0000000000..1793af2ba2 --- /dev/null +++ b/projects/clr/hipamd/.clang-format @@ -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 diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index baf3b49df1..1573ddee5c 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -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 diff --git a/projects/clr/hipamd/docs/markdown/hip_programming_guide.md b/projects/clr/hipamd/docs/markdown/hip_programming_guide.md index 5d0c1f2497..9313eb22e1 100644 --- a/projects/clr/hipamd/docs/markdown/hip_programming_guide.md +++ b/projects/clr/hipamd/docs/markdown/hip_programming_guide.md @@ -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. diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 9d0757f83a..7f159572d7 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -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 diff --git a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp index 02e2f1e524..f7de214f10 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp @@ -69,18 +69,16 @@ namespace hip_impl } }; - using RAII_global = std::unique_ptr; - const std::unordered_map< hsa_agent_t, std::vector>& executables(); const std::unordered_map< std::uintptr_t, std::vector>>& functions(); const std::unordered_map& function_names(); - std::unordered_map& globals(); + std::unordered_map& globals(); hsa_executable_t load_executable( const std::string& file, hsa_executable_t executable, hsa_agent_t agent); -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index dbd6d8b300..902e3620fa 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -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()); } diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 77526cf9ac..ea6462caf4 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -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); diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 45a44b3666..d173a2f295 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -568,7 +568,6 @@ hipError_t hipModuleGetTexRef( const auto it = globals().find(name); if (it == globals().end()) return ihipLogStatus(hipErrorInvalidValue); - *texRef = static_cast(it->second.get()); - + *texRef = reinterpret_cast(it->second); return ihipLogStatus(hipSuccess); } diff --git a/projects/clr/hipamd/src/program_state.cpp b/projects/clr/hipamd/src/program_state.cpp index e867887da2..aef92d463c 100644 --- a/projects/clr/hipamd/src/program_state.cpp +++ b/projects/clr/hipamd/src/program_state.cpp @@ -169,7 +169,7 @@ namespace lock_guard 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(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& globals() + unordered_map& globals() { - static unordered_map r; + static unordered_map 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. \ No newline at end of file + + // 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. diff --git a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp index c279658b5a..78f8de48a1 100644 --- a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp @@ -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 */