Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[HotFix] Adapt to changes to HIPRTC in current Staging/Mainline (future 6.1 RC) that disallow the use of <limits> in kernels #2650

Closed
wants to merge 5 commits into from
Closed

[HotFix] Adapt to changes to HIPRTC in current Staging/Mainline (future 6.1 RC) that disallow the use of <limits> in kernels #2650

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
ab17c2c
[HotFix] Customized limits header only applies to a range of compiler…
junliume Jan 2, 2024
860ab56
fix clang format issue
junliume Jan 2, 2024
210aeaf
Make workaround conditions consistent
junliume Jan 2, 2024
b3372a9
update WA info for lifecycle tracking
junliume Jan 3, 2024
7b4a17a
Add missing numeric_limits<int>
junliume Jan 3, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 8 additions & 3 deletions src/comgr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <limits> 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)
Expand Down Expand Up @@ -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 <limits> 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))
Expand Down
9 changes: 7 additions & 2 deletions src/kernels/miopen_cstdint.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <cstdint> // int8_t, int16_t
Expand Down
13 changes: 11 additions & 2 deletions src/kernels/miopen_limits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -41,6 +41,15 @@ namespace std {
template <typename T>
class numeric_limits;

template <>
Copy link
Collaborator Author

@junliume junliume Jan 3, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@atamazov could you take a look? test_reduce_test seems to be needing this in the docker where I am testing. Also I may wonder if there could be other data types missing?

class numeric_limits<int>
{
public:
static constexpr __device__ int max() noexcept { return __INT_MAX__; }

static constexpr __device__ int min() noexcept { return -__INT_MAX__ - 1; }
};

template <>
class numeric_limits<float>
{
Expand Down
7 changes: 5 additions & 2 deletions src/kernels/miopen_type_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@atamazov reviewers be noted that this line excludes the assumed (fake) version.


namespace std {

Expand Down Expand Up @@ -76,7 +79,7 @@ struct remove_cv
typedef typename remove_volatile<typename remove_const<T>::type>::type type;
};

#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL
#if WORKAROUND_USE_MORE_CUSTOM_TYPE_TRAITS
template <class T, T v>
struct integral_constant
{
Expand Down Expand Up @@ -105,7 +108,7 @@ using enable_if = __hip_internal::enable_if<B, T>;

template <bool B, typename T = void>
using enable_if_t = typename __hip_internal::enable_if<B, T>::type;
#endif
#endif // WORKAROUND_USE_MORE_CUSTOM_TYPE_TRAITS

template <class T>
struct is_pointer_helper : false_type
Expand Down
6 changes: 4 additions & 2 deletions test/handle_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <miopen/config.h>
#include <miopen/handle.hpp>
Expand Down