From 4b4eb5de805535e21ae3c01ac306e6fdef4cd175 Mon Sep 17 00:00:00 2001 From: RahulC <168588218+rahulc-gh@users.noreply.github.com> Date: Thu, 5 Mar 2026 08:29:46 -0800 Subject: [PATCH] Revert "Support fp8 types in hiprtc (#2605)" This reverts commit a1cf15ea9aeef472984e33e222d4057cbaf2f7b0. --- .../include/hip/amd_detail/amd_hip_fp8.h | 4 +- .../hip/amd_detail/amd_hip_mx_common.h | 2 - .../include/hip/amd_detail/amd_hip_ocp_fp.hpp | 15 +- .../hip/amd_detail/amd_hip_ocp_types.h | 24 +- projects/clr/hipamd/src/hiprtc/CMakeLists.txt | 4 - .../clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake | 1 + .../catch/unit/rtc/hiprtc_fp16_HeaderTst.cc | 310 +++++++++--------- .../hip-tests/catch/unit/rtc/hiprtc_fp8.cc | 132 -------- 8 files changed, 174 insertions(+), 318 deletions(-) delete mode 100644 projects/hip-tests/catch/unit/rtc/hiprtc_fp8.cc diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp8.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp8.h index c7320a710f4..fb4a97ca994 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp8.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp8.h @@ -64,9 +64,10 @@ #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 #include @@ -75,7 +76,6 @@ #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 diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_mx_common.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_mx_common.h index 970d9315883..7762c6e6935 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_mx_common.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_mx_common.h @@ -32,11 +32,9 @@ SOFTWARE. #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, diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp index d150e9249b7..6de994d293e 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp @@ -24,13 +24,14 @@ THE SOFTWARE. #pragma once -#if !defined(__HIPCC_RTC__) -#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 +#if !defined(__HIPCC_RTC__) +#include +#include +#include +#include +#include #include #include @@ -40,6 +41,8 @@ static_assert(sizeof(uint32_t) * CHAR_BIT == 32); static_assert(sizeof(uint64_t) * CHAR_BIT == 64); #endif // !defined(__HIPCC_RTC__) +#include // Host Conversion + // HW Detection #if defined(__gfx950__) #define HIP_ENABLE_GFX950_OCP_BUILTINS 1 diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h index ba5ded0c651..73773fee132 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h @@ -24,21 +24,11 @@ THE SOFTWARE. #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, ""); @@ -54,17 +44,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 __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 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; #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; 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; @@ -75,7 +65,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 __hip_uint32_t __attribute__((ext_vector_type(6))) __amd_fp6x32_storage_t; +typedef 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 */ diff --git a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt index 99989681a94..675d587f6ce 100644 --- a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt +++ b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt @@ -179,10 +179,6 @@ ${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. diff --git a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake index 7b7aefb9c08..5bb125d0093 100644 --- a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake +++ b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -129,3 +129,4 @@ __hipRTC_header:\n\ __hipRTC_header_size:\n\ .long __hipRTC_header_size - __hipRTC_header\n") endmacro(generate_hiprtc_mcin) + diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc index 8bbe26d44fe..5b2e0e82f8a 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc @@ -50,48 +50,48 @@ __global__ void fp16(float *res) { z = x - y; res[11] = __heq(z, 8); z = x * y; res[12] = __heq(z, 20); z = x / y; res[13] = __heq(z, 5); - x = 2.2; y = 2.22; res[14] = __heq(__half(x == y), 0); - x = 2.2; y = 2.2; res[15] = __heq(__half(x != y), 0); - x = 2.2; y = 2.22; res[16] = __heq(__half(x < y), 1); - x = 2.2; y = 2.3; res[17] = __heq(__half(x > y), 0); - x = 2.2; y = 2.3; res[18] = __heq(__half(x <= y), 1); - x = 2.2; y = 2.2; res[19] = __heq(__half(x >= y), 1); + x = 2.2; y = 2.22; res[14] = __heq((x == y), 0); + x = 2.2; y = 2.2; res[15] = __heq((x != y), 0); + x = 2.2; y = 2.22; res[16] = __heq((x < y), 1); + x = 2.2; y = 2.3; res[17] = __heq((x > y), 0); + x = 2.2; y = 2.3; res[18] = __heq((x <= y), 1); + x = 2.2; y = 2.2; res[19] = __heq((x >= y), 1); __half2 d = __half2{5, 8}, e = __half2{2, 10}, f = 0, g = __half2{3, 4}; - g += 2; res[20] = __heq(__half(__half2{5, 6} == g), 1); - g -= 2; res[21] = __heq(__half(__half2{3, 4} == g), 1); - g *= 2; res[22] = __heq(__half(__half2{6, 8} == g), 1); - g /= 2; res[23] = __heq(__half(__half2{3, 4} == g), 1); - g++; res[24] = __heq(__half(__half2{4, 5} == g), 1); - g--; res[25] = __heq(__half(__half2{3, 4} == g), 1); - ++g; res[26] = __heq(__half(__half2{4, 5} == g), 1); - --g; res[27] = __heq(__half(__half2{3, 4} == g), 1); - f = d + __half2{1, 1}; res[28] = __heq(__half(__half2{6, 9} == f), 1); - f = d - __half2{1, 1}; res[29] = __heq(__half(__half2{4, 7} == f), 1); - f = d + e; res[30] = __heq(__half(__half2{7, 18} == f), 1); - f = d - e; res[31] = __heq(__half(__half2{3, -2} == f), 1); - f = d * e; res[32] = __heq(__half(__half2{10, 80} == f), 1); - f = d / e; res[33] = __heq(__half(__half2{2.5, 0.8} == f), 1); - d = __half2{5, 8}; e = __half2{5.1, 7.9}; res[34] = __heq(__half(d == e), 0); - d = __half2{5, 8}; e = __half2{5.1, 8}; res[35] = __heq(__half(d != e), 1); - d = __half2{4, 7}; e = __half2{5, 8}; res[36] = __heq(__half(d < e), 1); - d = __half2{3, 8}; e = __half2{2, 8}; res[37] = __heq(__half(d > e), 0); - d = __half2{2, 8}; e = __half2{2, 10}; res[38] = __heq(__half(d <= e), 1); - d = __half2{5, 8}; e = __half2{2, 10}; res[39] = __heq(__half(d >= e), 1); + g += 2; res[20] = __heq((__half2{5, 6} == g), 1); + g -= 2; res[21] = __heq((__half2{3, 4} == g), 1); + g *= 2; res[22] = __heq((__half2{6, 8} == g), 1); + g /= 2; res[23] = __heq((__half2{3, 4} == g), 1); + g++; res[24] = __heq((__half2{4, 5} == g), 1); + g--; res[25] = __heq((__half2{3, 4} == g), 1); + ++g; res[26] = __heq((__half2{4, 5} == g), 1); + --g; res[27] = __heq((__half2{3, 4} == g), 1); + f = d + __half2{1, 1}; res[28] = __heq((__half2{6, 9} == f), 1); + f = d - __half2{1, 1}; res[29] = __heq((__half2{4, 7} == f), 1); + f = d + e; res[30] = __heq((__half2{7, 18} == f), 1); + f = d - e; res[31] = __heq((__half2{3, -2} == f), 1); + f = d * e; res[32] = __heq((__half2{10, 80} == f), 1); + f = d / e; res[33] = __heq((__half2{2.5, 0.8} == f), 1); + d = __half2{5, 8}; e = __half2{5.1, 7.9}; res[34] = __heq((d == e), 0); + d = __half2{5, 8}; e = __half2{5.1, 8}; res[35] = __heq((d != e), 1); + d = __half2{4, 7}; e = __half2{5, 8}; res[36] = __heq((d < e), 1); + d = __half2{3, 8}; e = __half2{2, 8}; res[37] = __heq((d > e), 0); + d = __half2{2, 8}; e = __half2{2, 10}; res[38] = __heq((d <= e), 1); + d = __half2{5, 8}; e = __half2{2, 10}; res[39] = __heq((d >= e), 1); - res[40] = __heq(__half(make_half2(2, 2) == __half2{2, 2}), 1); + res[40] = __heq((make_half2(2, 2) == __half2{2, 2}), 1); res[41] = __heq(__low2half(make_half2(2, 2)), 2); res[42] = __heq(__high2half(make_half2(3, 3)), 3); - res[43] = __heq(__half(__half2half2(3) == __half2{3, 3}), 1); - res[44] = __heq(__half(__halves2half2(3, 4) == __half2{3, 4}), 1); - res[45] = __heq(__half(__low2half2(__half2{2, 3}) == __half2{2, 2}), 1); - res[46] = __heq(__half(__high2half2(__half2{4, 5}) == __half2{5, 5}), 1); - res[47] = __heq(__half(__lows2half2(__half2{2, 4}, __half2{3, 5}) == __half2{2, 3}), 1); - res[48] = __heq(__half(__highs2half2(__half2{2, 4}, __half2{3, 5}) == __half2{4, 5}), 1); - res[49] = __heq(__half(__lowhigh2highlow(__half2{2, 3}) == __half2{3, 2}), 1); + res[43] = __heq((__half2half2(3) == __half2{3, 3}), 1); + res[44] = __heq((__halves2half2(3, 4) == __half2{3, 4}), 1); + res[45] = __heq((__low2half2(__half2{2, 3}) == __half2{2, 2}), 1); + res[46] = __heq((__high2half2(__half2{4, 5}) == __half2{5, 5}), 1); + res[47] = __heq((__lows2half2(__half2{2, 4}, __half2{3, 5}) == __half2{2, 3}), 1); + res[48] = __heq((__highs2half2(__half2{2, 4}, __half2{3, 5}) == __half2{4, 5}), 1); + res[49] = __heq((__lowhigh2highlow(__half2{2, 3}) == __half2{3, 2}), 1); - res[50] = __half_as_short(3) == 16896; - res[51] = __half_as_ushort(11) == 18816; + res[50] = __heq(__half_as_short(3), 16896); + res[51] = __heq(__half_as_ushort(11), 18816); res[52] = __heq(__short_as_half(16896), 3); res[53] = __heq(__ushort_as_half(18816), 11); @@ -100,57 +100,57 @@ __global__ void fp16(float *res) { res[56] = __heq(__float2half_rz(3.1234), 3.123047); res[57] = __heq(__float2half_rd(3.1234), 3.123047); res[58] = __heq(__float2half_ru(3.1234), 3.125); - res[59] = __heq(__half(__float2half2_rn(3) == __half2{3, 3}), 1); - res[60] = __heq(__half(__floats2half2_rn(3, 2) == __half2{3, 2}), 1); - res[61] = __heq(__half(__float22half2_rn(make_float2(3, 4)) == __half2{3, 4}), 1); - res[62] = __heq(__half2float(3.3), __half(3.3)); - res[63] = __heq(__low2float(__half2{3, 4}), __half(3)); - res[64] = __heq(__high2float(__half2{3, 4}), __half(4)); - res[65] = __heq(__half(__half22float2(__half2{3, 4}) == make_float2(3, 4)), 1); - res[66] = __half2int_rn(1.1234) == 1; - res[67] = __half2int_rz(1.1234) == 1; - res[68] = __half2int_rd(1.1234) == 1; - res[69] = __half2int_ru(1.1234) == 1; + 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); + res[63] = __heq(__low2float(__half2{3, 4}), 3); + res[64] = __heq(__high2float(__half2{3, 4}), 4); + res[65] = __heq((__half22float2(__half2{3, 4}) == make_float2(3, 4)), 1); + res[66] = __heq(__half2int_rn(1.1234), 1); + res[67] = __heq(__half2int_rz(1.1234), 1); + res[68] = __heq(__half2int_rd(1.1234), 1); + res[69] = __heq(__half2int_ru(1.1234), 1); res[70] = __heq(__int2half_rn(2), 2); res[71] = __heq(__int2half_rz(2), 2); res[72] = __heq(__int2half_rd(2), 2); res[73] = __heq(__int2half_ru(2), 2); - res[74] = __half2short_rn(1.1234) == 1; - res[75] = __half2short_rz(1.1234) == 1; - res[76] = __half2short_rd(1.1234) == 1; - res[77] = __half2short_ru(1.1234) == 1; + res[74] = __heq(__half2short_rn(1.1234), 1); + res[75] = __heq(__half2short_rz(1.1234), 1); + res[76] = __heq(__half2short_rd(1.1234), 1); + res[77] = __heq(__half2short_ru(1.1234), 1); res[78] = __heq(__short2half_rn(2), 2); res[79] = __heq(__short2half_rz(2), 2); res[80] = __heq(__short2half_rd(2), 2); res[81] = __heq(__short2half_ru(2), 2); - res[82] = __half2ll_rn(1.1234) == 1; - res[83] = __half2ll_rz(1.1234) == 1; - res[84] = __half2ll_rd(1.1234) == 1; - res[85] = __half2ll_ru(1.1234) == 1; + res[82] = __heq(__half2ll_rn(1.1234), 1); + res[83] = __heq(__half2ll_rz(1.1234), 1); + res[84] = __heq(__half2ll_rd(1.1234), 1); + res[85] = __heq(__half2ll_ru(1.1234), 1); res[86] = __heq(__ll2half_rn(2), 2); res[87] = __heq(__ll2half_rz(2), 2); res[88] = __heq(__ll2half_rd(2), 2); res[89] = __heq(__ll2half_ru(2), 2); - res[90] = __half2uint_rn(1.1234) == 1; - res[91] = __half2uint_rz(1.1234) == 1; - res[92] = __half2uint_rd(1.1234) == 1; - res[93] = __half2uint_ru(1.1234) == 1; + res[90] = __heq(__half2uint_rn(1.1234), 1); + res[91] = __heq(__half2uint_rz(1.1234), 1); + res[92] = __heq(__half2uint_rd(1.1234), 1); + res[93] = __heq(__half2uint_ru(1.1234), 1); res[94] = __heq(__uint2half_rn(2), 2); res[95] = __heq(__uint2half_rz(2), 2); res[96] = __heq(__uint2half_rd(2), 2); res[97] = __heq(__uint2half_ru(2), 2); - res[98] = __half2ushort_rn(1.1234) == 1; - res[99] = __half2ushort_rz(1.1234) == 1; - res[100] = __half2ushort_rd(1.1234) == 1; - res[101] = __half2ushort_ru(1.1234) == 1; + res[98] = __heq(__half2ushort_rn(1.1234), 1); + res[99] = __heq(__half2ushort_rz(1.1234), 1); + res[100] = __heq(__half2ushort_rd(1.1234), 1); + res[101] = __heq(__half2ushort_ru(1.1234), 1); res[102] = __heq(__ushort2half_rn(2), 2); res[103] = __heq(__ushort2half_rz(2), 2); res[104] = __heq(__ushort2half_rd(2), 2); res[105] = __heq(__ushort2half_ru(2), 2); - res[106] = __half2ull_rn(1.1234), 1; - res[107] = __half2ull_rz(1.1234), 1; - res[108] = __half2ull_rd(1.1234) == 1; - res[109] = __half2ull_ru(1.1234) == 1; + res[106] = __heq(__half2ull_rn(1.1234), 1); + res[107] = __heq(__half2ull_rz(1.1234), 1); + res[108] = __heq(__half2ull_rd(1.1234), 1); + res[109] = __heq(__half2ull_ru(1.1234), 1); res[110] = __heq(__ull2half_rn(2), 2); res[111] = __heq(__ull2half_rz(2), 2); res[112] = __heq(__ull2half_rd(2), 2); @@ -163,48 +163,48 @@ __global__ void fp16(float *res) { res[117] = __heq(__ldcs(&b), a); __half2 m, n; m = n = __half2{5, 7}; - res[118] = __heq(__half(__ldg(&m) == n), 1); - res[119] = __heq(__half(__ldcg(&m) == n), 1); - res[120] = __heq(__half(__ldca(&m) == n), 1); - res[121] = __heq(__half(__ldcs(&m) == n), 1); + res[118] = __heq((__ldg(&m) == n), 1); + res[119] = __heq((__ldcg(&m) == n), 1); + res[120] = __heq((__ldca(&m) == n), 1); + res[121] = __heq((__ldcs(&m) == n), 1); - a = 2.22; b = 2.22; res[122] = __heq(a, b) == 1; - a = 2.2; b = 2.22; res[123] = __hne(a, b) == 1; - a = 2.2; b = 2.201; res[124] = __hle(a, b) == 1; - a = 2.21; b = 2.201; res[125] = __hge(a, b) == 1; - a = 2.2; b = 2.201; res[126] = __hlt(a, b) == 1; - a = 2.2; b = 2.21; res[127] = __hgt(a, b) == 0; - a = 2.21; b = 2.201; res[128] = __hequ(a, b) == 0; - a = 2.201; b = 2.201; res[129] = __hneu(a, b) == 0; - a = 2.201; b = 2.21; res[130] = __hleu(a, b) == 1; - a = 2.21; b = 2.22; res[131] = __hgeu(a, b) == 0; - a = 2.201; b = 2.201; res[132] = __hltu(a, b) == 0; - a = 2.21; b = 2.201; res[133] = __hgtu(a, b) == 1; + a = 2.22; b = 2.22; res[122] = (__heq(a, b) == 1); + a = 2.2; b = 2.22; res[123] = __heq(__hne(a, b), 1); + a = 2.2; b = 2.201; res[124] = __heq(__hle(a, b), 1); + a = 2.21; b = 2.201; res[125] = __heq(__hge(a, b), 1); + a = 2.2; b = 2.201; res[126] = __heq(__hlt(a, b), 1); + a = 2.2; b = 2.21; res[127] = __heq(__hgt(a, b), 0); + a = 2.21; b = 2.201; res[128] = __heq(__hequ(a, b), 0); + a = 2.201; b = 2.201; res[129] = __heq(__hneu(a, b), 0); + a = 2.201; b = 2.21; res[130] = __heq(__hleu(a, b), 1); + a = 2.21; b = 2.22; res[131] = __heq(__hgeu(a, b), 0); + a = 2.201; b = 2.201; res[132] = __heq(__hltu(a, b), 0); + a = 2.21; b = 2.201; res[133] = __heq(__hgtu(a, b), 1); - res[134] = __heq(__half(__heq2(__half2{1, 2}, __half2{1, 2}) == __half2{1, 1}), 1); - res[135] = __heq(__half(__hne2(__half2{1, 2}, __half2{2, 2}) == __half2{1, 0}), 1); - res[136] = __heq(__half(__hle2(__half2{1, 2}, __half2{1, 1}) == __half2{1, 0}), 1); - res[137] = __heq(__half(__hge2(__half2{1, 2}, __half2{1, 1}) == __half2{1, 1}), 1); - res[138] = __heq(__half(__hlt2(__half2{1, 2}, __half2{2, 2}) == __half2{1, 0}), 1); - res[139] = __heq(__half(__hgt2(__half2{1, 2}, __half2{1, 1}) == __half2{0, 1}), 1); - res[140] = __heq(__half(__hequ2(__half2{2, 3}, __half2{2, 2}) == __half2{1, 0}), 1); - res[141] = __heq(__half(__hneu2(__half2{1, 2}, __half2{1, 4}) == __half2{0, 1}), 1); - res[142] = __heq(__half(__hleu2(__half2{2, 3}, __half2{2, 2}) == __half2{1, 0}), 1); - res[143] = __heq(__half(__hgeu2(__half2{2, 3}, __half2{3, 2}) == __half2{0, 1}), 1); - res[144] = __heq(__half(__hltu2(__half2{2, 3}, __half2{3, 2}) == __half2{1, 0}), 1); - res[145] = __heq(__half(__hgtu2(__half2{2, 3}, __half2{3, 2}) == __half2{0, 1}), 1); - res[146] = __hbeq2(__half2{3, 3}, __half2{3, -3}) == 0; - res[147] = __hbne2(__half2{4, 3}, __half2{3, 3}) == 0; - res[148] = __hble2(__half2{3, 3}, __half2{3, 4}) == 1; - res[149] = __hbge2(__half2{2, 4}, __half2{3, 3}) == 0; - res[150] = __hblt2(__half2{2, 4}, __half2{2, 4}) == 0; - res[151] = __hbgt2(__half2{2, 4}, __half2{2, 3}) == 0; - res[152] = __hbequ2(__half2{3, 2}, __half2{3, 3}) == 0; - res[153] = __hbneu2(__half2{4, 3}, __half2{4, 2}) == 0; - res[154] = __hbleu2(__half2{2, 3}, __half2{2, 3}) == 1; - res[155] = __hbgeu2(__half2{3, 3}, __half2{2, 3}) == 1; - res[156] = __hbltu2(__half2{2, 3}, __half2{2, 2}) == 0; - res[157] = __hbgtu2(__half2{3, 3}, __half2{2, 3}) == 0; + res[134] = __heq((__heq2(__half2{1, 2}, __half2{1, 2}) == __half2{1, 1}), 1); + res[135] = __heq((__hne2(__half2{1, 2}, __half2{2, 2}) == __half2{1, 0}), 1); + res[136] = __heq((__hle2(__half2{1, 2}, __half2{1, 1}) == __half2{1, 0}), 1); + res[137] = __heq((__hge2(__half2{1, 2}, __half2{1, 1}) == __half2{1, 1}), 1); + res[138] = __heq((__hlt2(__half2{1, 2}, __half2{2, 2}) == __half2{1, 0}), 1); + res[139] = __heq((__hgt2(__half2{1, 2}, __half2{1, 1}) == __half2{0, 1}), 1); + res[140] = __heq((__hequ2(__half2{2, 3}, __half2{2, 2}) == __half2{1, 0}), 1); + res[141] = __heq((__hneu2(__half2{1, 2}, __half2{1, 4}) == __half2{0, 1}), 1); + res[142] = __heq((__hleu2(__half2{2, 3}, __half2{2, 2}) == __half2{1, 0}), 1); + res[143] = __heq((__hgeu2(__half2{2, 3}, __half2{3, 2}) == __half2{0, 1}), 1); + res[144] = __heq((__hltu2(__half2{2, 3}, __half2{3, 2}) == __half2{1, 0}), 1); + res[145] = __heq((__hgtu2(__half2{2, 3}, __half2{3, 2}) == __half2{0, 1}), 1); + res[146] = __heq(__hbeq2(__half2{3, 3}, __half2{3, -3}), 0); + res[147] = __heq(__hbne2(__half2{4, 3}, __half2{3, 3}), 0); + res[148] = __heq(__hble2(__half2{3, 3}, __half2{3, 4}), 1); + res[149] = __heq(__hbge2(__half2{2, 4}, __half2{3, 3}), 0); + res[150] = __heq(__hblt2(__half2{2, 4}, __half2{2, 4}), 0); + res[151] = __heq(__hbgt2(__half2{2, 4}, __half2{2, 3}), 0); + res[152] = __heq(__hbequ2(__half2{3, 2}, __half2{3, 3}), 0); + res[153] = __heq(__hbneu2(__half2{4, 3}, __half2{4, 2}), 0); + res[154] = __heq(__hbleu2(__half2{2, 3}, __half2{2, 3}), 1); + res[155] = __heq(__hbgeu2(__half2{3, 3}, __half2{2, 3}), 1); + res[156] = __heq(__hbltu2(__half2{2, 3}, __half2{2, 2}), 0); + res[157] = __heq(__hbgtu2(__half2{3, 3}, __half2{2, 3}), 0); a = 2.2; b = 2.22; res[158] = __heq(__hmax(a, b), b); a = 2.2; b = 2.202; res[159] = __heq(__hmax_nan(a, b), b); @@ -224,55 +224,55 @@ __global__ void fp16(float *res) { a = 2.2; b = -1; res[173] = __heq(__hfma_sat(a, b, a), 0); a = 4.2; b = 2; res[174] = __heq(__hdiv(a, b), 2.1); - res[175] = __hbeq2(__habs2(__half2{-1, -4}), __half2{1, 4}) == 1; - res[176] = __hbeq2(__hadd2(__half2{1, 4}, __half2{2, -5}), __half2{3, -1}) == 1; - res[177] = __hbeq2(__hsub2(__half2{1, 4}, __half2{2, -2}), __half2{-1, 6}) == 1; - res[178] = __hbeq2(__hmul2(__half2{1, 3}, __half2{5, -2}), __half2{5, -6}) == 1; - res[179] = __hbeq2(__hadd2_sat(__half2{1, 3}, __half2{2, -5}), __half2{1, 0}) == 1; - res[180] = __hbeq2(__hsub2_sat(__half2{2, 3}, __half2{2, -2}), __half2{0, 1}) == 1; - res[181] = __hbeq2(__hmul2_sat(__half2{1, 3}, __half2{5, -2}), __half2{1, 0}) == 1; - res[182] = __hbeq2(__hfma2(__half2{1, 3}, __half2{5, -2}, __half2{-5, 8}), __half2{0, 2}) == 1; - res[183] = __hbeq2(__hfma2_sat(__half2{1, 3}, __half2{5, -2}, __half2{-5, 8}), __half2{0, 1}) == 1; - res[184] = __hbeq2(__h2div(__half2{1, 3}, __half2{5, -2}), __half2{0.2, -1.5}) == 1; - res[185] = amd_mixed_dot(__half2{1, 3}, __half2{3, 3}, 2, 1) == 14; + res[175] = __heq(__hbeq2(__habs2(__half2{-1, -4}), __half2{1, 4}), 1); + res[176] = __heq(__hbeq2(__hadd2(__half2{1, 4}, __half2{2, -5}), __half2{3, -1}), 1); + res[177] = __heq(__hbeq2(__hsub2(__half2{1, 4}, __half2{2, -2}), __half2{-1, 6}), 1); + res[178] = __heq(__hbeq2(__hmul2(__half2{1, 3}, __half2{5, -2}), __half2{5, -6}), 1); + res[179] = __heq(__hbeq2(__hadd2_sat(__half2{1, 3}, __half2{2, -5}), __half2{1, 0}), 1); + res[180] = __heq(__hbeq2(__hsub2_sat(__half2{2, 3}, __half2{2, -2}), __half2{0, 1}), 1); + res[181] = __heq(__hbeq2(__hmul2_sat(__half2{1, 3}, __half2{5, -2}), __half2{1, 0}), 1); + res[182] = __heq(__hbeq2(__hfma2(__half2{1, 3}, __half2{5, -2}, __half2{-5, 8}), __half2{0, 2}), 1); + res[183] = __heq(__hbeq2(__hfma2_sat(__half2{1, 3}, __half2{5, -2}, __half2{-5, 8}), __half2{0, 1}), 1); + res[184] = __heq(__hbeq2(__h2div(__half2{1, 3}, __half2{5, -2}), __half2{0.2, -1.5}), 1); + res[185] = __heq(amd_mixed_dot(__half2{1, 3}, __half2{3, 3}, 2, 1), 14); - res[186] = __heq(htrunc(__half(2.8)), 2); - res[187] = __heq(hceil(__half(2.8)), 3); - res[188] = __heq(hfloor(__half(2.8)), 2); - res[189] = __heq(hrint(__half(2.8)), 3); - res[190] = __heq(hsin(__half(0)), 0); - res[191] = __heq(hcos(__half(0)), 1); - res[192] = __heq(hexp(__half(2)), 7.390625); - res[193] = __heq(hexp2(__half(2)), 4); - res[194] = __heq(hexp10(__half(2)), 100); - res[195] = __heq(hlog(__half(7.390625)), 2); - res[196] = __heq(hlog2(__half(4)), 2); - res[197] = __heq(hlog10(__half(100)), 2); - res[198] = __heq(hrcp(__half(4)), 0.25); - res[199] = __heq(hrsqrt(__half(0.25)), 2); - res[200] = __heq(hsqrt(__half(1.21)), 1.1); - 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); + res[187] = __heq(hceil(2.8), 3); + res[188] = __heq(hfloor(2.8), 2); + 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); + res[193] = __heq(hexp2(2), 4); + res[194] = __heq(hexp10(2), 100); + res[195] = __heq(hlog(7.390625), 2); + res[196] = __heq(hlog2(4), 2); + res[197] = __heq(hlog10(100), 2); + res[198] = __heq(hrcp(4), 0.25); + res[199] = __heq(hrsqrt(0.25), 2); + 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); - res[204] = __hbeq2(h2trunc(__half2{3.4, 5.2}), __half2{3, 5}) == 1; - res[205] = __hbeq2(h2ceil(__half2{3.4, 5.2}), __half2{4, 6}) == 1; - res[206] = __hbeq2(h2floor(__half2{3.4, 5.2}), __half2{3, 5}) == 1; - res[207] = __hbeq2(h2rint(__half2{3.4, 5.2}), __half2{3, 5}) == 1; - res[208] = __hbeq2(h2sin(__half2{0, 0}), __half2{0, 0}) == 1; - res[209] = __hbeq2(h2cos(__half2{0, 0}), __half2{1, 1}) == 1; - res[210] = __hbeq2(h2exp(__half2{2, 0}), __half2{7.390625, 1}) == 1; - res[211] = __hbeq2(h2exp2(__half2{3, 2}), __half2{8, 4}) == 1; - res[212] = __hbeq2(h2exp10(__half2{2, 3}), __half2{100, 1000}) == 1; - res[213] = __hbeq2(h2log(__half2{2.718750, 1}), __half2{1, 0}) == 1; - res[214] = __hbeq2(h2log2(__half2{8, 16}), __half2{3, 4}) == 1; - res[215] = __hbeq2(h2log10(__half2{1000, 100}), __half2{3, 2}) == 1; - res[216] = __hbeq2(h2rcp(__half2{4, 5}), __half2{0.25, 0.2}) == 1; - res[217] = __hbeq2(h2rsqrt(__half2{100, 25}), __half2{0.1, 0.2}) == 1; - res[218] = __hbeq2(h2sqrt(__half2{100, 25}), __half2{10, 5}) == 1; - res[219] = __hbeq2(__hisinf2(__half2{100, 0}), __half2{0, 0}) == 1; - res[220] = __hbeq2(__hisnan2(__half2{100, 25}), __half2{0, 0}) == 1; - res[221] = __hbeq2(__hneg2(__half2{2.1, -25}), __half2{-2.1, 25}) == 1; + res[204] = __heq(__hbeq2(h2trunc(__half2{3.4, 5.2}), __half2{3, 5}), 1); + res[205] = __heq(__hbeq2(h2ceil(__half2{3.4, 5.2}), __half2{4, 6}), 1); + res[206] = __heq(__hbeq2(h2floor(__half2{3.4, 5.2}), __half2{3, 5}), 1); + res[207] = __heq(__hbeq2(h2rint(__half2{3.4, 5.2}), __half2{3, 5}), 1); + res[208] = __heq(__hbeq2(h2sin(__half2{0, 0}), __half2{0, 0}), 1); + res[209] = __heq(__hbeq2(h2cos(__half2{0, 0}), __half2{1, 1}), 1); + res[210] = __heq(__hbeq2(h2exp(__half2{2, 0}), __half2{7.390625, 1}), 1); + res[211] = __heq(__hbeq2(h2exp2(__half2{3, 2}), __half2{8, 4}), 1); + res[212] = __heq(__hbeq2(h2exp10(__half2{2, 3}), __half2{100, 1000}), 1); + res[213] = __heq(__hbeq2(h2log(__half2{2.718750, 1}), __half2{1, 0}), 1); + res[214] = __heq(__hbeq2(h2log2(__half2{8, 16}), __half2{3, 4}), 1); + res[215] = __heq(__hbeq2(h2log10(__half2{1000, 100}), __half2{3, 2}), 1); + res[216] = __heq(__hbeq2(h2rcp(__half2{4, 5}), __half2{0.25, 0.2}), 1); + res[217] = __heq(__hbeq2(h2rsqrt(__half2{100, 25}), __half2{0.1, 0.2}), 1); + res[218] = __heq(__hbeq2(h2sqrt(__half2{100, 25}), __half2{10, 5}), 1); + res[219] = __heq(__hbeq2(__hisinf2(__half2{100, 0}), __half2{0, 0}), 1); + res[220] = __heq(__hbeq2(__hisnan2(__half2{100, 25}), __half2{0, 0}), 1); + res[221] = __heq(__hbeq2(__hneg2(__half2{2.1, -25}), __half2{-2.1, 25}), 1); } )"}; @@ -317,7 +317,7 @@ TEST_CASE("Unit_Rtc_fp16_header") { HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); std::string log(logSize, '\0'); HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); - std::cout << log << std::endl; + WARN(log); REQUIRE(false); } size_t codeSize; diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_fp8.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_fp8.cc deleted file mode 100644 index 7b648f6b40a..00000000000 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_fp8.cc +++ /dev/null @@ -1,132 +0,0 @@ -/* -Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - - -#include - -#include -#include - -#include -#include - -TEST_CASE("Unit_hiprtc_fp8_simple") { - constexpr const char* source = R"( -extern "C" __global__ void float_to_fp8_to_float(float* out, float* in, bool e4m3, size_t size) { - size_t i = threadIdx.x; - if (i < size) { - if (e4m3) { - __hip_fp8_e4m3 tmp = in[i]; - out[i] = tmp; - } else { - __hip_fp8_e5m2 tmp = in[i]; - out[i] = tmp; - } - } -} -)"; - - hiprtcProgram prog; - HIPRTC_CHECK(hiprtcCreateProgram(&prog, source, "fp8.cu", 0, nullptr, nullptr)); - hipDeviceProp_t props; - int device = 0; - HIP_CHECK(hipGetDeviceProperties(&props, device)); -#ifdef __HIP_PLATFORM_AMD__ - std::string sarg = std::string("--offload-arch=") + props.gcnArchName; -#else - std::string sarg = std::string("--fmad=false"); -#endif - const char* options[] = {sarg.c_str()}; - hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; - size_t logSize; - HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); - if (logSize) { - std::string log(logSize, '\0'); - HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); - std::cout << log << '\n'; - } - REQUIRE(compileResult == HIPRTC_SUCCESS); - - size_t codeSize{}; - HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); - - std::vector code(codeSize); - HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); - - HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); - - constexpr size_t size = 10; - - float *d_in, *d_out; - HIP_CHECK(hipMalloc(&d_in, size * sizeof(float))); - HIP_CHECK(hipMalloc(&d_out, size * sizeof(float))); - - hipModule_t module; - hipFunction_t kernel; - HIP_CHECK(hipModuleLoadData(&module, code.data())); - HIP_CHECK(hipModuleGetFunction(&kernel, module, "float_to_fp8_to_float")); - - std::vector in(size, 0.0f); - for (size_t i = 0; i < size; i++) { - in[i] = -5.0f + i; - } - - HIP_CHECK(hipMemcpy(d_in, in.data(), size * sizeof(float), hipMemcpyHostToDevice)); - HIP_CHECK(hipMemset(d_out, 0, size * sizeof(float))); - - struct { - float* out; - float* in; - bool e4m3; - size_t size; - } args{d_out, d_in, true, size}; - - auto arg_size = sizeof(args); - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &arg_size, - HIP_LAUNCH_PARAM_END}; - - HIP_CHECK(hipModuleLaunchKernel(kernel, 1, 1, 1, size, 1, 1, 0, nullptr, nullptr, config)); - - std::vector out(size, 0.0f); - HIP_CHECK(hipMemcpy(out.data(), d_out, size * sizeof(float), hipMemcpyDeviceToHost)); - - for (size_t i = 0; i < size; i++) { - __hip_fp8_e4m3 tmp = in[i]; - float cpu_out = tmp; - INFO("Index: " << i << " in: " << in[i] << " GPU: " << out[i] << " cpu: " << cpu_out); - REQUIRE(cpu_out == out[i]); - } - - args.e4m3 = false; - HIP_CHECK(hipMemset(d_out, 0, size * sizeof(float))); - HIP_CHECK(hipModuleLaunchKernel(kernel, 1, 1, 1, size, 1, 1, 0, nullptr, nullptr, config)); - - HIP_CHECK(hipMemcpy(out.data(), d_out, size * sizeof(float), hipMemcpyDeviceToHost)); - for (size_t i = 0; i < size; i++) { - __hip_fp8_e5m2 tmp = in[i]; - float cpu_out = tmp; - INFO("Index: " << i << " in: " << in[i] << " GPU: " << out[i] << " cpu: " << cpu_out); - REQUIRE(cpu_out == out[i]); - } - - HIP_CHECK(hipFree(d_in)); - HIP_CHECK(hipFree(d_out)); - - HIP_CHECK(hipModuleUnload(module)); -}