diff --git a/src/comgr.cpp b/src/comgr.cpp index 971a60391e..3df75213a8 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -60,6 +60,12 @@ /// With base driver 5.11.32 the errors disappear. /// More info at https://github.com/ROCm/MIOpen/issues/1257. #define WORKAROUND_ISSUE_1257 (HIP_PACKAGE_VERSION_FLAT >= 4003021331ULL) +/// For now, use only standard to avoid possibility of +/// correctnes or performance regressions. +/// \todo Test and enable "custom" local implementation. +/// \todo: exact version from which this WA is needed should be updated in release +/// https://github.com/ROCm/MIOpen/issues/2651 +#define WORKAROUND_DONT_USE_CUSTOM_LIMITS (HIP_PACKAGE_VERSION_FLAT < 6001024000ULL) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMGR_LOG_CALLS) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMGR_LOG_SOURCE_NAMES) @@ -1302,10 +1308,9 @@ void BuildHip(const std::string& name, #endif opts.push_back("-DHIP_PACKAGE_VERSION_FLAT=" + std::to_string(HIP_PACKAGE_VERSION_FLAT)); opts.push_back("-DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS"); - /// For now, use only standard to avoid possibility of - /// correctnes or performance regressions. - /// \todo Test and enable "custom" local implementation. +#if WORKAROUND_DONT_USE_CUSTOM_LIMITS opts.push_back("-DWORKAROUND_DONT_USE_CUSTOM_LIMITS=1"); +#endif #if WORKAROUND_ISSUE_1431 if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) && !miopen::comgr::IsWave64Enforced(opts)) diff --git a/src/kernels/miopen_cstdint.hpp b/src/kernels/miopen_cstdint.hpp index 57d0d088b1..985aad1c80 100644 --- a/src/kernels/miopen_cstdint.hpp +++ b/src/kernels/miopen_cstdint.hpp @@ -26,15 +26,20 @@ #pragma once #ifdef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS +/// \todo: exact version from which this WA is needed should be updated in release +/// https://github.com/ROCm/MIOpen/issues/2651 +#define WORKAROUND_USE_MORE_CUSTOM_CSTDINT (HIP_PACKAGE_VERSION_FLAT >= 6001024000ULL) + typedef signed char int8_t; typedef unsigned char uint8_t; typedef signed short int16_t; typedef unsigned short uint16_t; -#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL + +#if WORKAROUND_USE_MORE_CUSTOM_CSTDINT typedef signed int int32_t; typedef unsigned int uint32_t; typedef __hip_internal::uint64_t uint64_t; -#endif +#endif // WORKAROUND_USE_MORE_CUSTOM_CSTDINT #else #include // int8_t, int16_t diff --git a/src/kernels/miopen_limits.hpp b/src/kernels/miopen_limits.hpp index 2a8f5e6178..d3c89ca0f9 100644 --- a/src/kernels/miopen_limits.hpp +++ b/src/kernels/miopen_limits.hpp @@ -25,8 +25,8 @@ *******************************************************************************/ #pragma once -#ifndef WORKAROUND_DO_NOT_USE_CUSTOM_LIMITS -#define WORKAROUND_DO_NOT_USE_CUSTOM_LIMITS 0 +#ifndef WORKAROUND_DONT_USE_CUSTOM_LIMITS +#define WORKAROUND_DONT_USE_CUSTOM_LIMITS 0 #endif #if defined(MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS) && !WORKAROUND_DONT_USE_CUSTOM_LIMITS @@ -41,6 +41,15 @@ namespace std { template class numeric_limits; +template <> +class numeric_limits +{ +public: + static constexpr __device__ int max() noexcept { return __INT_MAX__; } + + static constexpr __device__ int min() noexcept { return -__INT_MAX__ - 1; } +}; + template <> class numeric_limits { diff --git a/src/kernels/miopen_type_traits.hpp b/src/kernels/miopen_type_traits.hpp index 352010e3ac..8ef6a8c80b 100644 --- a/src/kernels/miopen_type_traits.hpp +++ b/src/kernels/miopen_type_traits.hpp @@ -26,6 +26,9 @@ #pragma once #ifdef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS +/// \todo: exact version from which this WA is needed should be updated in release +/// https://github.com/ROCm/MIOpen/issues/2651 +#define WORKAROUND_USE_MORE_CUSTOM_TYPE_TRAITS (HIP_PACKAGE_VERSION_FLAT > 6001024000ULL) namespace std { @@ -76,7 +79,7 @@ struct remove_cv typedef typename remove_volatile::type>::type type; }; -#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL +#if WORKAROUND_USE_MORE_CUSTOM_TYPE_TRAITS template struct integral_constant { @@ -105,7 +108,7 @@ using enable_if = __hip_internal::enable_if; template using enable_if_t = typename __hip_internal::enable_if::type; -#endif +#endif // WORKAROUND_USE_MORE_CUSTOM_TYPE_TRAITS template struct is_pointer_helper : false_type diff --git a/test/handle_test.cpp b/test/handle_test.cpp index ba143d7ad1..b726fd8377 100644 --- a/test/handle_test.cpp +++ b/test/handle_test.cpp @@ -29,12 +29,14 @@ #define WORKAROUND_SWDEV_257056_PCH_MISSING_MACROS 1 // https://gerrit-git.amd.com/c/compute/ec/clr/+/972441 -// "HIP_PACKAGE_VERSION_FLAT == 6001000000ULL" is for ROCm 6.1 RC where issue #2600 is not +// "HIP_PACKAGE_VERSION_FLAT == 6001024000ULL" is for ROCm 6.1 RC where issue #2600 is not // yet fixed in the compiler. In order to test such release candidates, we have to // override HIP version to 6.1.0. +/// \todo: exact version from which this WA is needed should be updated in release +/// https://github.com/ROCm/MIOpen/issues/2651 #define WORKAROUND_ISSUE_2600 \ ((HIP_PACKAGE_VERSION_FLAT >= 6000000000ULL && HIP_PACKAGE_VERSION_FLAT <= 6000999999ULL) || \ - HIP_PACKAGE_VERSION_FLAT == 6001000000ULL) + HIP_PACKAGE_VERSION_FLAT == 6001024000ULL) #include #include