diff options
Diffstat (limited to 'dev-util/hip/files')
| -rw-r--r-- | dev-util/hip/files/hip-5.7.1-disable-stack-protector.patch | 13 | ||||
| -rw-r--r-- | dev-util/hip/files/hip-5.7.1-exec-stack.patch | 31 | ||||
| -rw-r--r-- | dev-util/hip/files/hip-5.7.1-fix-unaligned-access.patch | 67 | ||||
| -rw-r--r-- | dev-util/hip/files/hip-5.7.1-no_asan_doc.patch | 17 |
4 files changed, 128 insertions, 0 deletions
diff --git a/dev-util/hip/files/hip-5.7.1-disable-stack-protector.patch b/dev-util/hip/files/hip-5.7.1-disable-stack-protector.patch new file mode 100644 index 000000000000..49d944ef4bc9 --- /dev/null +++ b/dev-util/hip/files/hip-5.7.1-disable-stack-protector.patch @@ -0,0 +1,13 @@ +Disable stack-protector (which is enabled by default gentoo-hardened) in device code. +This is not required after https://github.com/llvm/llvm-project/pull/70799, but helps with older clang. +Upstream bug: https://github.com/ROCm-Developer-Tools/clr/issues/21 +--- a/hipamd/src/hiprtc/hiprtcInternal.cpp ++++ b/hipamd/src/hiprtc/hiprtcInternal.cpp +@@ -144,6 +144,7 @@ RTCCompileProgram::RTCCompileProgram(std::string name_) : RTCProgram(name_), fgp + compile_options_.push_back("-nogpuinc"); + compile_options_.push_back("-Wno-gnu-line-marker"); + compile_options_.push_back("-Wno-missing-prototypes"); ++ compile_options_.push_back("-fno-stack-protector"); + #ifdef _WIN32 + compile_options_.push_back("-target"); + compile_options_.push_back("x86_64-pc-windows-msvc"); diff --git a/dev-util/hip/files/hip-5.7.1-exec-stack.patch b/dev-util/hip/files/hip-5.7.1-exec-stack.patch new file mode 100644 index 000000000000..744ca7b953c3 --- /dev/null +++ b/dev-util/hip/files/hip-5.7.1-exec-stack.patch @@ -0,0 +1,31 @@ +Upstream bug: https://github.com/ROCm-Developer-Tools/clr/issues/22 +--- a/hipamd/src/hip_embed_pch.sh ++++ b/hipamd/src/hip_embed_pch.sh +@@ -178,6 +178,7 @@ EOF + + echo "// Automatically generated script for HIP RTC." > $mcinFile + if [[ $isWindows -eq 0 ]]; then ++ echo " .section .note.GNU-stack,"",%progbits" >> $mcinFile + echo " .type __hipRTC_header,@object" >> $mcinFile + echo " .type __hipRTC_header_size,@object" >> $mcinFile + fi +--- a/hipamd/src/hiprtc/cmake/HIPRTC.cmake ++++ b/hipamd/src/hiprtc/cmake/HIPRTC.cmake +@@ -98,6 +98,7 @@ macro(generate_hiprtc_mcin HiprtcMcin HiprtcPreprocessedInput) + set(HIPRTC_TYPE_LINUX_ONLY "") + else() + set(HIPRTC_TYPE_LINUX_ONLY ++ " .section .note.GNU-stack,\"\",%progbits\n" + " .type __hipRTC_header,@object\n" + " .type __hipRTC_header_size,@object") + endif() +--- a/hipamd/src/hip_embed_pch.sh ++++ b/hipamd/src/hip_embed_pch.sh +@@ -111,6 +111,7 @@ cat >$tmp/hip_pch.h <<EOF + EOF + + cat >$tmp/hip_pch.mcin <<EOF ++ .section .note.GNU-stack,"",%progbits + .type __hip_pch_wave32,@object + .section .hip_pch_wave32,"aMS",@progbits,1 + .data diff --git a/dev-util/hip/files/hip-5.7.1-fix-unaligned-access.patch b/dev-util/hip/files/hip-5.7.1-fix-unaligned-access.patch new file mode 100644 index 000000000000..ae2092f6e1ac --- /dev/null +++ b/dev-util/hip/files/hip-5.7.1-fix-unaligned-access.patch @@ -0,0 +1,67 @@ +Fix SIGSEGV when compiled with avx-512 instructions. + +Due to unaligned allocations, library crashes in +nontemporalMemcpy in _mm512_stream_si512 (which requires +64-aligned allocations, but used to copy default-aligned objects). + +Without this patch hipamd causes random crashes in hipMemcpy* callers +(tensile, rocBLAS, miopen, rocThrust, etc.). + +Bug: https://bugs.gentoo.org/915969 +Bug report in upstream: https://github.com/ROCm-Developer-Tools/clr/issues/18 +--- a/rocclr/device/rocm/rocvirtual.cpp ++++ b/rocclr/device/rocm/rocvirtual.cpp +@@ -2790,44 +2790,6 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize) + return true; + } + +-// ================================================================================================ +-__attribute__((optimize("unroll-all-loops"), always_inline)) +-static inline void nontemporalMemcpy(void* __restrict dst, const void* __restrict src, +- uint16_t size) { +- #if defined(__AVX512F__) +- for (auto i = 0u; i != size / sizeof(__m512i); ++i) { +- _mm512_stream_si512(reinterpret_cast<__m512i* __restrict&>(dst)++, +- *reinterpret_cast<const __m512i* __restrict&>(src)++); +- } +- size = size % sizeof(__m512i); +- #endif +- +- #if defined(__AVX__) +- for (auto i = 0u; i != size / sizeof(__m256i); ++i) { +- _mm256_stream_si256(reinterpret_cast<__m256i* __restrict&>(dst)++, +- *reinterpret_cast<const __m256i* __restrict&>(src)++); +- } +- size = size % sizeof(__m256i); +- #endif +- +- for (auto i = 0u; i != size / sizeof(__m128i); ++i) { +- _mm_stream_si128(reinterpret_cast<__m128i* __restrict&>(dst)++, +- *(reinterpret_cast<const __m128i* __restrict&>(src)++)); +- } +- size = size % sizeof(__m128i); +- +- for (auto i = 0u; i != size / sizeof(long long); ++i) { +- _mm_stream_si64(reinterpret_cast<long long* __restrict&>(dst)++, +- *reinterpret_cast<const long long* __restrict&>(src)++); +- } +- size = size % sizeof(long long); +- +- for (auto i = 0u; i != size / sizeof(int); ++i) { +- _mm_stream_si32(reinterpret_cast<int* __restrict&>(dst)++, +- *reinterpret_cast<const int* __restrict&>(src)++); +- } +-} +- + // ================================================================================================ + bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, + const amd::Kernel& kernel, const_address parameters, void* eventHandle, +@@ -3096,7 +3058,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, + argBuffer = reinterpret_cast<address>(allocKernArg(gpuKernel.KernargSegmentByteSize(), + gpuKernel.KernargSegmentAlignment())); + // Load all kernel arguments +- nontemporalMemcpy(argBuffer, parameters, ++ memcpy(argBuffer, parameters, + std::min(gpuKernel.KernargSegmentByteSize(), + signature.paramsSize())); + } diff --git a/dev-util/hip/files/hip-5.7.1-no_asan_doc.patch b/dev-util/hip/files/hip-5.7.1-no_asan_doc.patch new file mode 100644 index 000000000000..40de416334be --- /dev/null +++ b/dev-util/hip/files/hip-5.7.1-no_asan_doc.patch @@ -0,0 +1,17 @@ +Donot install -asan documents. + +Reference: +https://github.com/RadeonOpenCompute/llvm-project/commit/e782e09f7b113a0f896c6cec7240d411aca1d73f +https://github.com/RadeonOpenCompute/ROCm-CompilerSupport/issues/61 + +--- a/hipamd/packaging/CMakeLists.txt 2023-11-18 23:25:45.000000000 +0800 ++++ b/hipamd/packaging/CMakeLists.txt 2023-11-18 23:27:27.230354665 +0800 +@@ -36,8 +36,6 @@ + ###Set License#### + set(CPACK_RESOURCE_FILE_LICENSE ${hip_SOURCE_DIR}/LICENSE.txt) + install(FILES ${CPACK_RESOURCE_FILE_LICENSE} DESTINATION ${CMAKE_INSTALL_DOCDIR} COMPONENT binary) +-# install license file in share/doc/hip-asan folder +-install(FILES ${CPACK_RESOURCE_FILE_LICENSE} DESTINATION ${CMAKE_INSTALL_DOCDIR}-asan COMPONENT asan) + set(CPACK_RPM_PACKAGE_LICENSE "MIT") + #Begin binary files install + if(HIP_PLATFORM STREQUAL "amd" ) |