Revert "Support fp8 types in hiprtc" by rahulc-gh · Pull Request #3792 · ROCm/rocm-systems

rahulc-gh · 2026-03-05T16:30:57Z

Reverts #2605

rocm-systems bump PR ROCm/TheRock#3754, failing fft/miopen/rccl/rocprofiler-compute tests needs this revert to unblock

Following tickets have been raised by Wang, Yanyao
rocfft/hipfft:
ROCm/TheRock#3771
hipdnn-samples:
ROCm/TheRock#3772
rccl:
ROCm/TheRock#3774
miopen:
ROCm/TheRock#3776
miopenprovider:
ROCm/TheRock#3777
rocprofiler_compute:
ROCm/TheRock#3778

This reverts commit a1cf15e.

Copilot

Pull request overview

Reverts the prior change that enabled FP8 types in HIPRTC to unblock downstream ROCm component test failures.

Changes:

Removes the HIPRTC FP8 unit test.
Stops packaging FP8/BF16/OCP-related headers into the HIPRTC header bundle.
Adjusts the FP16 HIPRTC header test to avoid problematic constructs and to standardize result comparisons/logging.

Reviewed changes

Copilot reviewed 8 out of 8 changed files in this pull request and generated 10 comments.

Show a summary per file

File	Description
projects/hip-tests/catch/unit/rtc/hiprtc_fp8.cc	Removes the FP8 HIPRTC unit test that depends on the reverted functionality.
projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc	Refactors comparisons/validation patterns in the FP16 RTC header test and switches failure logging to Catch2 diagnostics.
projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake	No functional change (trailing newline).
projects/clr/hipamd/src/hiprtc/CMakeLists.txt	Removes FP8/OCP-related headers from the HIPRTC packaged header set.
projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h	Removes RTC-specific macro path and switches storage typedefs to standard fixed-width integer types.
projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp	Reorders/changes includes; makes host conversion header include unconditional.
projects/clr/hipamd/include/hip/amd_detail/amd_hip_mx_common.h	Removes RTC guard around internal includes.
projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp8.h	Adjusts include ordering/guards around BF16/MX common dependencies.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Copilot · 2026-03-05T17:47:48Z

+typedef uint8_t __amd_fp8_storage_t;
+typedef uint16_t __amd_fp8x2_storage_t;
+typedef uint8_t __amd_fp4x2_storage_t;
+typedef uint32_t __amd_fp4x8_storage_t;
 typedef __bf16 __amd_bf16_storage_t;
 typedef _Float16 __amd_fp16_storage_t;
-typedef __hip_int8_t __amd_scale_t;
+typedef int8_t __amd_scale_t;


This header now uses fixed-width integer types (uint8_t/uint16_t/uint32_t/int8_t) but does not include a header that guarantees these typedefs are available. Add an include for (or <stdint.h>) near the top of the file to avoid compilation failures depending on include order.

Copilot · 2026-03-05T17:47:48Z

 #if defined(__clang__) && (__clang_major__ > 17) && defined(__HIP__)
 typedef unsigned int __attribute__((ext_vector_type(2))) __amd_uintx2_storage_t;
-typedef __hip_uint8_t __attribute__((ext_vector_type(8))) __amd_fp8x8_storage_t;
+typedef uint8_t __attribute__((ext_vector_type(8))) __amd_fp8x8_storage_t;


This header now uses fixed-width integer types (uint8_t/uint16_t/uint32_t/int8_t) but does not include a header that guarantees these typedefs are available. Add an include for (or <stdint.h>) near the top of the file to avoid compilation failures depending on include order.

Copilot · 2026-03-05T17:47:48Z

 typedef _Float16 __attribute__((ext_vector_type(8))) __amd_fp16x8_storage_t;
 typedef _Float16 __attribute__((ext_vector_type(32))) __amd_fp16x32_storage_t;
-typedef __hip_uint32_t __attribute__((ext_vector_type(6))) __amd_fp6x32_storage_t;
+typedef uint32_t __attribute__((ext_vector_type(6))) __amd_fp6x32_storage_t;


This header now uses fixed-width integer types (uint8_t/uint16_t/uint32_t/int8_t) but does not include a header that guarantees these typedefs are available. Add an include for (or <stdint.h>) near the top of the file to avoid compilation failures depending on include order.

Copilot · 2026-03-05T17:47:49Z

 #define __OCP_FP_HOST__ __host__
 #define __OCP_FP_DEVICE__ __device__
 #define __OCP_FP_HOST_DEVICE__ __OCP_FP_HOST__ __OCP_FP_DEVICE__
 #define __OCP_FP_DEVICE_STATIC__ __OCP_FP_DEVICE__ static __inline__ __attribute__((always_inline))
 #define __OCP_FP_HOST_DEVICE_STATIC__ __OCP_FP_HOST_DEVICE__ static


The previous version included host_defines.h (in the non-RTC path) to ensure host/device are defined. With that include removed, this header may fail when included from translation units that don't already define these macros. Consider including the appropriate HIP defines header (e.g., hip/amd_detail/host_defines.h) before using host/device, or provide safe fallbacks when they are not defined.

Copilot · 2026-03-05T17:47:49Z

+#if !defined(__HIPCC_RTC__)
+#include <hip/amd_detail/amd_hip_common.h>
+#include <hip/amd_detail/host_defines.h>
+#include <hip/amd_detail/amd_hip_ocp_types.h>
+#include <hip/hip_fp16.h>
+#include <hip/hip_bf16.h>


The comment indicates amd_hip_ocp_host.hpp is for host conversions, but it is now included unconditionally (including under HIPCC_RTC). If this header is host-only (or pulls in host-only dependencies), it should remain inside the !defined(HIPCC_RTC) block, or be made explicitly RTC-safe to keep include intent clear and avoid RTC build issues.

Copilot · 2026-03-05T17:47:49Z

 #endif  // !defined(__HIPCC_RTC__)

+#include <hip/amd_detail/amd_hip_ocp_host.hpp>  // Host Conversion
+
 // HW Detection
 #if defined(__gfx950__)


The comment indicates amd_hip_ocp_host.hpp is for host conversions, but it is now included unconditionally (including under HIPCC_RTC). If this header is host-only (or pulls in host-only dependencies), it should remain inside the !defined(HIPCC_RTC) block, or be made explicitly RTC-safe to keep include intent clear and avoid RTC build issues.

Suggested change

#endif // !defined(__HIPCC_RTC__)

#include <hip/amd_detail/amd_hip_ocp_host.hpp> // Host Conversion

// HW Detection

#if defined(__gfx950__)

#include <hip/amd_detail/amd_hip_ocp_host.hpp> // Host Conversion

#endif // !defined(__HIPCC_RTC__)

// HW Detection

#if defined(__gfx950__)

#if defined(__gfx950__)

Copilot · 2026-03-05T17:47:50Z

+  res[59] = __heq((__float2half2_rn(3) == __half2{3, 3}), 1);
+  res[60] = __heq((__floats2half2_rn(3, 2) == __half2{3, 2}), 1);
+  res[61] = __heq((__float22half2_rn(make_float2(3, 4)) == __half2{3, 4}), 1);
+  res[62] = __heq(__half2float(3.3), 3.3);


These changes introduce several double-precision literals (e.g., 3.3, 2.8, 1.25) into device code where float/half behavior is being tested. To avoid unintended overload selection and implicit double-to-float/half conversions (and related warnings), prefer using explicit float literals (e.g., 3.3f, 2.8f, 1.25f) or explicit casts consistent with the function under test.

Copilot · 2026-03-05T17:47:50Z

-  res[201] = __hisinf(__half(1)) == 0;
-  res[202] = __hisnan(__half(1)) == 0;
-  res[203] = __heq(__hneg(__half(1.25)), -1.25);
+  res[186] = __heq(htrunc(2.8), 2);


These changes introduce several double-precision literals (e.g., 3.3, 2.8, 1.25) into device code where float/half behavior is being tested. To avoid unintended overload selection and implicit double-to-float/half conversions (and related warnings), prefer using explicit float literals (e.g., 3.3f, 2.8f, 1.25f) or explicit casts consistent with the function under test.

Copilot · 2026-03-05T17:47:50Z

+  res[189] = __heq(hrint(2.8), 3);
+  res[190] = __heq(hsin(0), 0);
+  res[191] = __heq(hcos(0), 1);
+  res[192] = __heq(hexp(2), 7.390625);


These changes introduce several double-precision literals (e.g., 3.3, 2.8, 1.25) into device code where float/half behavior is being tested. To avoid unintended overload selection and implicit double-to-float/half conversions (and related warnings), prefer using explicit float literals (e.g., 3.3f, 2.8f, 1.25f) or explicit casts consistent with the function under test.

Copilot · 2026-03-05T17:47:50Z

+  res[200] = __heq(hsqrt(1.21), 1.1);
+  res[201] = __heq(__hisinf(1), 0);
+  res[202] = __heq(__hisnan(1), 0);
+  res[203] = __heq(__hneg(1.25), -1.25);


These changes introduce several double-precision literals (e.g., 3.3, 2.8, 1.25) into device code where float/half behavior is being tested. To avoid unintended overload selection and implicit double-to-float/half conversions (and related warnings), prefer using explicit float literals (e.g., 3.3f, 2.8f, 1.25f) or explicit casts consistent with the function under test.

Suggested change

res[203] = __heq(__hneg(1.25), -1.25);

res[203] = __heq(__hneg(1.25f), -1.25f);

radhaksri

Approved

This reverts commit a1cf15ea9aeef472984e33e222d4057cbaf2f7b0. [rocm-systems] ROCm/rocm-systems#3792 (commit be0675a)

This reverts commit be0675a.

This reverts commit be0675a (#3792) This change was earlier reverted due to issues seen in MIOpen via the rocm-systems bump PR ROCm/TheRock#3754. Now, MIOpen has fixed the issues using https://github.com/ROCm/rocm-libraries/pull/5235/changes. So reverting the revert to bring back the support for fp8 types in hiprtc. ## Motivation Users should be able to include fp8 header in hiprtc code. ## Technical Details Fix some include issues in fp8 header. ## JIRA ID ROCM-864 ## Test Plan Added an fp8 test for hiprtc, basically include it from HIP directory and use it in rtc code. ## Test Result The test basically hiprtc output of fp8 conversion round trip float to what we get on CPU. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Jatin Chaudhary <jatchaud@amd.com> Co-authored-by: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com>

Revert "Support fp8 types in hiprtc (#2605)"

4b4eb5d

This reverts commit a1cf15e.

Copilot AI review requested due to automatic review settings March 5, 2026 16:30

rahulc-gh requested review from a team as code owners March 5, 2026 16:30

Copilot AI reviewed Mar 5, 2026

View reviewed changes

systems-assistant Bot added the organization: ROCm label Mar 5, 2026

rahulc-gh requested review from chrispaquot, cjatin and radhaksri March 5, 2026 18:47

radhaksri approved these changes Mar 5, 2026

View reviewed changes

rahulc-gh merged commit be0675a into develop Mar 5, 2026
23 of 31 checks passed

rahulc-gh deleted the revert-2605-users/jatin/support_fp8_hiprtc branch March 5, 2026 22:38

systems-assistant Bot pushed a commit to ROCm/clr that referenced this pull request Mar 5, 2026

Revert "Support fp8 types in hiprtc (#2605)"

0a29b9a

This reverts commit a1cf15ea9aeef472984e33e222d4057cbaf2f7b0. [rocm-systems] ROCm/rocm-systems#3792 (commit be0675a)

systems-assistant Bot pushed a commit to ROCm/hip-tests that referenced this pull request Mar 5, 2026

Revert "Support fp8 types in hiprtc (#2605)"

acaeb7f

This reverts commit a1cf15ea9aeef472984e33e222d4057cbaf2f7b0. [rocm-systems] ROCm/rocm-systems#3792 (commit be0675a)

satyanveshd mentioned this pull request Apr 6, 2026

Revert "Revert "Support fp8 types in hiprtc (#2605)" (#3792)" #4727

Merged

1 task

cjatin pushed a commit that referenced this pull request Apr 7, 2026

Revert "Revert "Support fp8 types in hiprtc (#2605)" (#3792)"

1bbecc1

This reverts commit be0675a.

ScottTodd mentioned this pull request May 18, 2026

[Issue] rocfft and hipfft tests crash with SIGABRT on Linux gfx94X-dcgpu after rocm-systems bump ROCm/TheRock#3771

Closed

amd-chiranjeevi mentioned this pull request May 20, 2026

[Issue] hipdnn-samples: all 15 tests fail on Linux gfx94X-dcgpu — MIOpen "named symbol not found" after rocm-systems bump ROCm/TheRock#3772

Closed

Navigation Menu

Uh oh!

Revert "Support fp8 types in hiprtc" - #3792

Revert "Support fp8 types in hiprtc"#3792
rahulc-gh merged 1 commit into
developfrom
revert-2605-users/jatin/support_fp8_hiprtc

rahulc-gh commented Mar 5, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

Copilot AI Mar 5, 2026

Uh oh!

radhaksri left a comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

	res[203] = __heq(__hneg(1.25), -1.25);
	res[203] = __heq(__hneg(1.25f), -1.25f);

Uh oh!

Conversation

rahulc-gh commented Mar 5, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Mar 5, 2026

Choose a reason for hiding this comment

Uh oh!

radhaksri left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants