Skip to content

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

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

Merged
cjatin merged 4 commits into from
Apr 23, 2026
Merged

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

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
1bbecc1
Revert "Revert "Support fp8 types in hiprtc (#2605)" (#3792)"
satyanveshd Apr 6, 2026
5d83ddb
Address copilot comments
satyanveshd Apr 6, 2026
2d34c2e
add if/else for fp8 ocp/fnuz
cjatin Apr 7, 2026
f814fa2
Merge branch 'develop' into revert-3792-users/jatin/support_fp8_hiprtc
cjatin Apr 22, 2026
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
5 changes: 2 additions & 3 deletions projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp8.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,11 +45,9 @@
#endif
#endif

// Include it explicitly for HIPRTC
#if !defined(__HIPCC_RTC__)
#include "amd_hip_bf16.h"
#include "amd_hip_mx_common.h"

#if !defined(__HIPCC_RTC__)
#include <hip/amd_detail/amd_hip_common.h>
#include <climits>

Expand All @@ -58,6 +56,7 @@
#include "amd_hip_fp16.h" // __half_raw
#include "math_fwd.h" // ocml device functions
#include "hip_assert.h" // hip assertions

#define __HIP_SCHAR_MAX SCHAR_MAX
#define __HIP_SCHAR_MIN SCHAR_MIN
#define __HIP_UCHAR_MAX UCHAR_MAX
Expand Down
2 changes: 2 additions & 0 deletions projects/clr/hipamd/include/hip/amd_detail/amd_hip_mx_common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,9 +16,11 @@
#define HIP_ENABLE_HOST_OCP_CONVERSIONS 0
#endif

#if !defined(__HIPCC_RTC__)
#include "amd_hip_ocp_types.h"
#include "amd_hip_fp16.h"
#include "amd_hip_bf16.h"
#endif

enum hipRoundMode {
hipRoundNearest = 0,
Expand Down
15 changes: 6 additions & 9 deletions projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,14 +6,13 @@

#pragma once

#include <hip/amd_detail/amd_hip_bf16.h>

#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>
#include "amd_hip_common.h"
#include "host_defines.h"
#include "amd_hip_ocp_types.h"
#include "amd_hip_bf16.h"
#include "amd_hip_ocp_host.hpp"

#include <climits>
#include <cstdio>

Expand All @@ -23,8 +22,6 @@ static_assert(sizeof(uint32_t) * CHAR_BIT == 32);
static_assert(sizeof(uint64_t) * CHAR_BIT == 64);
#endif // !defined(__HIPCC_RTC__)

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

// HW Detection
#if defined(__gfx950__)
#define HIP_ENABLE_GFX950_OCP_BUILTINS 1
Expand Down
24 changes: 17 additions & 7 deletions projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,11 +6,21 @@

#pragma once

#if defined(__HIPCC_RTC__)
#define __OCP_FP_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
#else
#include "host_defines.h"

#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
#endif

static_assert(sizeof(unsigned int) == 4, "");
static_assert(sizeof(float) == 4, "");
Expand All @@ -26,17 +36,17 @@ static_assert(sizeof(_Float16) == 2, "");
// header which will act as a base abstraction, and will be maintained in the future, it makes sense
// to keep these vector types separate from existing implementations. We can add conversion
// functions in a different header using these functions.
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 __hip_uint8_t __amd_fp8_storage_t;
typedef __hip_uint16_t __amd_fp8x2_storage_t;
typedef __hip_uint8_t __amd_fp4x2_storage_t;
typedef __hip_uint32_t __amd_fp4x8_storage_t;
typedef __bf16 __amd_bf16_storage_t;
typedef _Float16 __amd_fp16_storage_t;
typedef int8_t __amd_scale_t;
typedef __hip_int8_t __amd_scale_t;

#if defined(__clang__) && (__clang_major__ > 17) && defined(__HIP__)
typedef unsigned int __attribute__((ext_vector_type(2))) __amd_uintx2_storage_t;
typedef uint8_t __attribute__((ext_vector_type(8))) __amd_fp8x8_storage_t;
typedef __hip_uint8_t __attribute__((ext_vector_type(8))) __amd_fp8x8_storage_t;
typedef __bf16 __attribute__((ext_vector_type(2))) __amd_bf16x2_storage_t;
typedef __bf16 __attribute__((ext_vector_type(8))) __amd_bf16x8_storage_t;
typedef __bf16 __attribute__((ext_vector_type(32))) __amd_bf16x32_storage_t;
Expand All @@ -47,7 +57,7 @@ typedef float __attribute__((ext_vector_type(32))) __amd_floatx32_storage_t;
typedef _Float16 __attribute__((ext_vector_type(2))) __amd_fp16x2_storage_t;
typedef _Float16 __attribute__((ext_vector_type(8))) __amd_fp16x8_storage_t;
typedef _Float16 __attribute__((ext_vector_type(32))) __amd_fp16x32_storage_t;
typedef uint32_t __attribute__((ext_vector_type(6))) __amd_fp6x32_storage_t;
typedef __hip_uint32_t __attribute__((ext_vector_type(6))) __amd_fp6x32_storage_t;
typedef short __attribute__((ext_vector_type(2))) __amd_shortx2_storage_t;
#elif defined(__GNUC__) && (__GNUC__ > 13)
/* GCC expects vector size in bytes */
Expand Down
4 changes: 4 additions & 0 deletions projects/clr/hipamd/src/hiprtc/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -167,6 +167,10 @@ ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/math_fwd.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_fp16_math_fwd.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_fp16.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_math_functions.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_bf16.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_ocp_types.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_mx_common.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_fp8.h
)

# Generate required HIPRTC files.
Expand Down
1 change: 0 additions & 1 deletion projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -112,4 +112,3 @@ __hipRTC_header:\n\
__hipRTC_header_size:\n\
.long __hipRTC_header_size - __hipRTC_header\n")
endmacro(generate_hiprtc_mcin)

1 change: 1 addition & 0 deletions projects/hip-tests/catch/unit/rtc/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@ set(AMD_TEST_SRC
hipRtcPtrdiff_t.cc
hiprtc_Bitcode_UndefinedFn.cc
hiprtc_bfloat16_HeaderTst.cc
hiprtc_fp8.cc
)

add_custom_target(copyRtcHeaders ALL
Expand Down
Loading
Loading