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

junliume · 2024-01-02T21:25:40Z

as commented below, with this PR, one has to use

add -DMIOPEN_OVERRIDE_HIP_VERSION_MINOR=1 -DMIOPEN_OVERRIDE_HIP_VERSION_PATCH=24000

to CMake options

Workaround lifecycle is tracked by #2651

…/hipRTC version

junliume · 2024-01-02T21:38:48Z

src/comgr.cpp

+/// Hence the workaround is applicable for a range of hipRTC versions
+/// and potentially in the future
+#define WORKAROUND_HIPRTC_LIMITS \
+    (HIP_PACKAGE_VERSION_FLAT >= 6000023464ULL && HIP_PACKAGE_VERSION_FLAT < 6000023494ULL)


For reviewers (@atamazov), these way points are the best we know at the moment, they may not be exactly precise but should work with official releases.

The HIP version in Staging is usually incorrect. Therefore, we should not rely on these numbers in the code (otherwise we will have more and more of mess in the source code). It's much better to replace the incorrect HIP version numbers with the correct expected ones (in this case 6.1.24xxx) and use the correct version.

@atamazov here is a summary of compiler versions and reported in the cmake output:

[Docker A] ROCm 6.0:

HIP version: 6.0.32830-d62f6a171

roc-6.0.0 23483 7208e8d15fbf218deb74483ea8c549c67ca4985e

-- Build with HIP 6.0.23494

[Docker B] The docker with all proposed hipRTC changes:

HIP version: 6.0.32830-dc7635e0d

roc-6.0.0 23483 7208e8d15fbf218deb74483ea8c549c67ca4985e

-- Build with HIP 6.0.23494

[Docker C] The latest mainline staging docker:

HIP version: 6.0.32830-606beddb1

amd-mainline-open 23512 8f0cb360b515076885176a3e5ba09dab6ecefa20

-- Build with HIP 6.0.23464

There are a few critical issues:

From Docker A and Docker B: hipRTC versioning is now decoupled from HIP package, so our code path control needs to accomodate.

From Docker C: the reported HIP version is incorrect

atamazov

Why: #2650 (comment)

⚠️ In order to use the library with current staging/mainline compiler, the users have to add -DMIOPEN_OVERRIDE_HIP_VERSION_MINOR=1 -DMIOPEN_OVERRIDE_HIP_VERSION_PATCH=24000 to CMake options. Please inform QA about this.

src/comgr.cpp

atamazov · 2024-01-02T22:43:09Z

@junliume Maybe let's rename PR to "[HotFix] Adapt to changes to HIPRTC in current Staging/Mainline (future 6.1 RC) that disallow the use of <limits> in kernels"?

junliume · 2024-01-02T23:34:17Z

src/kernels/miopen_type_traits.hpp

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

-#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL
+#if HIP_PACKAGE_VERSION_FLAT >= 6001024000ULL


@atamazov now with the header issue resolved, we found more issues with

/tmp/comgr-4e3a20/include/miopen_type_traits.hpp:81:8: error: redefinition of 'integral_constant' 81 | struct integral_constant | ^ /long_pathname_so_that_rpms_can_package_the_debug_info/src/external/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h:55:38: note: previous definition is here 55 | template <class _Tp, _Tp __v> struct integral_constant { | ^

I guess we have to identify another point where this workaround is needed, and not the same point with the limits header.

We can make it

#if HIP_PACKAGE_VERSION_FLAT > 6001024000ULL

but this is very fragile ...

junliume · 2024-01-02T23:52:14Z

src/kernels/miopen_cstdint.hpp

@@ -30,7 +30,7 @@ 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 HIP_PACKAGE_VERSION_FLAT >= 6001024000ULL


Similar to https://github.com/ROCm/MIOpen/pull/2650/files#r1439946058

I am afraid this is fragile and might or might not break ...

junliume · 2024-01-03T00:42:42Z

src/kernels/miopen_limits.hpp

@@ -41,6 +41,15 @@ namespace std {
 template <typename T>
 class numeric_limits;

+template <>


@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?

junliume · 2024-01-03T06:45:49Z

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)


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

junliume · 2024-01-04T06:43:16Z

replace with #2652

junliume added 2 commits January 2, 2024 13:25

[HotFix] Customized limits header only applies to a range of compiler…

ab17c2c

…/hipRTC version

fix clang format issue

860ab56

junliume commented Jan 2, 2024

View reviewed changes

junliume added non-miopen-bug urgency_blocker workaround value_unknown labels Jan 2, 2024

atamazov suggested changes Jan 2, 2024

View reviewed changes

src/comgr.cpp Outdated Show resolved Hide resolved

src/comgr.cpp Outdated Show resolved Hide resolved

Make workaround conditions consistent

210aeaf

junliume changed the title ~~[HotFix] Customized limits header only applies to a range of compiler/hipRTC version~~ [HotFix] Adapt to changes to HIPRTC in current Staging/Mainline (future 6.1 RC) that disallow the use of <limits> in kernels Jan 2, 2024

junliume commented Jan 2, 2024

View reviewed changes

junliume mentioned this pull request Jan 3, 2024

[HIP Package] HIP version from cmake is unreliable #2651

Open

junliume added 2 commits January 2, 2024 16:06

update WA info for lifecycle tracking

b3372a9

Add missing numeric_limits<int>

7b4a17a

junliume commented Jan 3, 2024

View reviewed changes

junliume requested a review from JehandadKhan January 3, 2024 06:46

CAHEK7 mentioned this pull request Jan 3, 2024

[ROCm 6.0.1] Automatically activate the new HIPRTC PCH adaptations starting from the 6.0.24000 version. Fix some build errors. #2644

Merged

atamazov mentioned this pull request Jan 4, 2024

[HOTFIX] Adapt to changes in HIP Mainline 417 (possibly future 6.1 RC) #2652

Merged

junliume closed this Jan 4, 2024

junliume deleted the fix_limits_apply branch February 6, 2024 20:10

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

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

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

junliume commented Jan 2, 2024 •

edited

Loading

junliume Jan 2, 2024

atamazov Jan 2, 2024

junliume Jan 2, 2024 •

edited

Loading

atamazov left a comment •

edited

Loading

atamazov commented Jan 2, 2024 •

edited

Loading

junliume Jan 2, 2024 •

edited

Loading

junliume Jan 2, 2024

junliume Jan 2, 2024

junliume Jan 3, 2024 •

edited

Loading

junliume Jan 3, 2024

junliume commented Jan 4, 2024

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

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

Conversation

junliume commented Jan 2, 2024 • edited Loading

junliume Jan 2, 2024

Choose a reason for hiding this comment

atamazov Jan 2, 2024

Choose a reason for hiding this comment

junliume Jan 2, 2024 • edited Loading

Choose a reason for hiding this comment

atamazov left a comment • edited Loading

Choose a reason for hiding this comment

atamazov commented Jan 2, 2024 • edited Loading

junliume Jan 2, 2024 • edited Loading

Choose a reason for hiding this comment

junliume Jan 2, 2024

Choose a reason for hiding this comment

junliume Jan 2, 2024

Choose a reason for hiding this comment

junliume Jan 3, 2024 • edited Loading

Choose a reason for hiding this comment

junliume Jan 3, 2024

Choose a reason for hiding this comment

junliume commented Jan 4, 2024

junliume commented Jan 2, 2024 •

edited

Loading

junliume Jan 2, 2024 •

edited

Loading

atamazov left a comment •

edited

Loading

atamazov commented Jan 2, 2024 •

edited

Loading

junliume Jan 2, 2024 •

edited

Loading

junliume Jan 3, 2024 •

edited

Loading