diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index fa58c3e45d..edac461eb0 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2023 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 @@ -24,13 +24,12 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H -#include - -#include "hip/amd_detail/host_defines.h" #if defined(__HIPCC_RTC__) #define __HOST_DEVICE__ __device__ #else #define __HOST_DEVICE__ __host__ __device__ + #include + #include "hip/amd_detail/host_defines.h" #include #if defined(__cplusplus) #include @@ -39,7 +38,7 @@ THE SOFTWARE. #endif #endif // !defined(__HIPCC_RTC__) -#if __HIP_CLANG_ONLY__ +#if defined(__clang__) && defined(__HIP__) typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2))); struct __half_raw { @@ -64,10 +63,13 @@ THE SOFTWARE. }; #if defined(__cplusplus) + #if !defined(__HIPCC_RTC__) #include "hip_fp16_math_fwd.h" #include "amd_hip_vector_types.h" #include "host_defines.h" - + #include "amd_device_functions.h" + #include "amd_warp_functions.h" + #endif namespace std { template<> struct is_floating_point<_Float16> : std::true_type {}; @@ -1471,7 +1473,7 @@ THE SOFTWARE. } // Math functions - #if __HIP_CLANG_ONLY__ + #if defined(__clang__) && defined(__HIP__) inline __device__ float amd_mixed_dot(__half2 a, __half2 b, float c, bool saturate) { @@ -1710,69 +1712,67 @@ THE SOFTWARE. using half = __half; using half2 = __half2; #endif - #if !defined(__HIPCC_RTC__) - #include "amd_device_functions.h" - #include "amd_warp_functions.h" - __device__ - inline - __half __shfl(__half var, int src_lane, int width = warpSize) { - union { int i; __half h; } tmp; tmp.h = var; - tmp.i = __shfl(tmp.i, src_lane, width); - return tmp.h; - } - __device__ - inline - __half2 __shfl(__half2 var, int src_lane, int width = warpSize) { - union { int i; __half2 h; } tmp; tmp.h = var; - tmp.i = __shfl(tmp.i, src_lane, width); - return tmp.h; - } - __device__ - inline - __half __shfl_up(__half var, unsigned int lane_delta, int width = warpSize) { - union { int i; __half h; } tmp; tmp.h = var; - tmp.i = __shfl_up(tmp.i, lane_delta, width); - return tmp.h; - } - __device__ - inline - __half2 __shfl_up(__half2 var, unsigned int lane_delta, int width = warpSize) { - union { int i; __half2 h; } tmp; tmp.h = var; - tmp.i = __shfl_up(tmp.i, lane_delta, width); - return tmp.h; - } - __device__ - inline - __half __shfl_down(__half var, unsigned int lane_delta, int width = warpSize) { - union { int i; __half h; } tmp; tmp.h = var; - tmp.i = __shfl_down(tmp.i, lane_delta, width); - return tmp.h; - } - __device__ - inline - __half2 __shfl_down(__half2 var, unsigned int lane_delta, int width = warpSize) { - union { int i; __half2 h; } tmp; tmp.h = var; - tmp.i = __shfl_down(tmp.i, lane_delta, width); - return tmp.h; - } - __device__ - inline - __half __shfl_xor(__half var, int lane_mask, int width = warpSize) { - union { int i; __half h; } tmp; tmp.h = var; - tmp.i = __shfl_xor(tmp.i, lane_mask, width); - return tmp.h; - } - __device__ - inline - __half2 __shfl_xor(__half2 var, int lane_mask, int width = warpSize) { - union { int i; __half2 h; } tmp; tmp.h = var; - tmp.i = __shfl_xor(tmp.i, lane_mask, width); - return tmp.h; - } - #endif // !defined(__HIPCC_RTC__) + __device__ + inline + __half __shfl(__half var, int src_lane, int width = warpSize) { + union { int i; __half h; } tmp; tmp.h = var; + tmp.i = __shfl(tmp.i, src_lane, width); + return tmp.h; + } + __device__ + inline + __half2 __shfl(__half2 var, int src_lane, int width = warpSize) { + union { int i; __half2 h; } tmp; tmp.h = var; + tmp.i = __shfl(tmp.i, src_lane, width); + return tmp.h; + } + __device__ + inline + __half __shfl_up(__half var, unsigned int lane_delta, int width = warpSize) { + union { int i; __half h; } tmp; tmp.h = var; + tmp.i = __shfl_up(tmp.i, lane_delta, width); + return tmp.h; + } + __device__ + inline + __half2 __shfl_up(__half2 var, unsigned int lane_delta, int width = warpSize) { + union { int i; __half2 h; } tmp; tmp.h = var; + tmp.i = __shfl_up(tmp.i, lane_delta, width); + return tmp.h; + } + __device__ + inline + __half __shfl_down(__half var, unsigned int lane_delta, int width = warpSize) { + union { int i; __half h; } tmp; tmp.h = var; + tmp.i = __shfl_down(tmp.i, lane_delta, width); + return tmp.h; + } + __device__ + inline + __half2 __shfl_down(__half2 var, unsigned int lane_delta, int width = warpSize) { + union { int i; __half2 h; } tmp; tmp.h = var; + tmp.i = __shfl_down(tmp.i, lane_delta, width); + return tmp.h; + } + __device__ + inline + __half __shfl_xor(__half var, int lane_mask, int width = warpSize) { + union { int i; __half h; } tmp; tmp.h = var; + tmp.i = __shfl_xor(tmp.i, lane_mask, width); + return tmp.h; + } + __device__ + inline + __half2 __shfl_xor(__half2 var, int lane_mask, int width = warpSize) { + union { int i; __half2 h; } tmp; tmp.h = var; + tmp.i = __shfl_xor(tmp.i, lane_mask, width); + return tmp.h; + } #endif // defined(__cplusplus) #elif defined(__GNUC__) - #include "hip_fp16_gcc.h" + #if !defined(__HIPCC_RTC__) + #include "hip_fp16_gcc.h" + #endif #endif // !defined(__clang__) && defined(__GNUC__) #endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h b/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h index caf6ec752c..a4ae27ca6f 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2023 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 @@ -25,7 +25,9 @@ THE SOFTWARE. // /* // Half Math Functions // */ +#if !defined(__HIPCC_RTC__) #include "host_defines.h" +#endif #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ extern "C" { @@ -55,7 +57,7 @@ extern "C" typedef _Float16 __2f16 __attribute__((ext_vector_type(2))); typedef short __2i16 __attribute__((ext_vector_type(2))); - #if __HIP_CLANG_ONLY__ + #if defined(__clang__) && defined(__HIP__) __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b, float c, bool s); #endif diff --git a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt index e2c409d511..90b7330e97 100644 --- a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt +++ b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt @@ -130,6 +130,8 @@ set(HIPRTC_GEN_MCIN "${HIPRTC_GEN_DIR}/hipRTC_header.mcin") set(HIPRTC_GEN_PREPROCESSED "${HIPRTC_GEN_DIR}/hipRTC") set(HIPRTC_GEN_OBJ "${HIPRTC_GEN_DIR}/hipRTC_header${CMAKE_CXX_OUTPUT_EXTENSION}") set(HIPRTC_WARP_FUNCS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_functions.h") +set(HIPRTC_FP16_MATH_FWD "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_fp16_math_fwd.h") +set(HIPRTC_FP16_FUNCS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_fp16.h") set(HIPRTC_COOP_GROUPS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_cooperative_groups.h") set(HIPRTC_COOP_GRPS_HELPER "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_cooperative_groups_helper.h") set(HIPRTC_UNSAFE_ATOMICS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_unsafe_atomics.h") @@ -145,7 +147,7 @@ generate_hiprtc_mcin("${HIPRTC_GEN_MCIN}" "${HIPRTC_GEN_PREPROCESSED}") add_custom_command( OUTPUT ${HIPRTC_GEN_PREPROCESSED} COMMAND $ -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++17 -nogpulib --hip-version=3.6 -isystem ${HIP_COMMON_INCLUDE_DIR} -isystem ${PROJECT_SOURCE_DIR}/include -isystem ${PROJECT_BINARY_DIR}/include -isystem ${CMAKE_CURRENT_SOURCE_DIR}/include --cuda-device-only -D__HIPCC_RTC__ -x hip ${HIPRTC_GEN_HEADER} -E -o ${HIPRTC_GEN_PREPROCESSED} - COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_WARP_HEADER_FILE=${HIPRTC_WARP_FUNCS} -DHIPRTC_COOP_HEADER_FILE=${HIPRTC_COOP_GROUPS} -DHIPRTC_COOP_HELPER_FILE=${HIPRTC_COOP_GRPS_HELPER} -DHIPRTC_UNSAFE_ATOMICS_FILE=${HIPRTC_UNSAFE_ATOMICS} -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE} + COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_WARP_HEADER_FILE=${HIPRTC_WARP_FUNCS} -DHIPRTC_COOP_HEADER_FILE=${HIPRTC_COOP_GROUPS} -DHIPRTC_COOP_HELPER_FILE=${HIPRTC_COOP_GRPS_HELPER} -DHIPRTC_UNSAFE_ATOMICS_FILE=${HIPRTC_UNSAFE_ATOMICS} -DHIPRTC_FP16_MATH_FWD_FILE=${HIPRTC_FP16_MATH_FWD} -DHIPRTC_FP16_HEADER_FILE=${HIPRTC_FP16_FUNCS} -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE} DEPENDS clang ${HIPRTC_GEN_HEADER}) add_custom_command( OUTPUT ${HIPRTC_GEN_OBJ} diff --git a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake index 521b02345d..730d013e28 100644 --- a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake +++ b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -51,7 +51,6 @@ function(get_hiprtc_macros HIPRTC_DEFINES) select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__) \n\ #pragma clang diagnostic pop\n\ #define HIP_INCLUDE_HIP_HIP_RUNTIME_H\n\ -#define HIP_INCLUDE_HIP_HIP_FP16_H\n\ #pragma clang diagnostic push\n\ #pragma clang diagnostic ignored \"-Wreserved-macro-identifier\"\n\ #define _HIP_BFLOAT16_H_\n\ @@ -75,6 +74,10 @@ if(HIPRTC_ADD_MACROS) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_COOP_HEADER}") FILE(READ "${HIPRTC_UNSAFE_ATOMICS_FILE}" HIPRTC_UNSAFE_ATOMICS) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_UNSAFE_ATOMICS}") + FILE(READ "${HIPRTC_FP16_MATH_FWD_FILE}" HIPRTC_FP16_MATH_FWD) + FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_FP16_MATH_FWD}") + FILE(READ "${HIPRTC_FP16_HEADER_FILE}" HIPRTC_FP16_HEADER) + FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_FP16_HEADER}") #pragma clang diagnostic pop endif() @@ -85,7 +88,6 @@ macro(generate_hiprtc_header HiprtcHeader) #define CHAR_BIT __CHAR_BIT__\n\ #define INT_MAX __INTMAX_MAX__\n\ #include \"hip/hip_runtime.h\"\n\ -#include \"hip/hip_fp16.h\"\n\ #include \"hip/hip_bfloat16.h\"\n\ #pragma pop_macro(\"CHAR_BIT\")\n\ #pragma pop_macro(\"INT_MAX\")")