AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-384714 - Add hipRTC support for amd_hip_fp16

Browse Source 
Change-Id: Ib6e2b46908d740909dfc15253965590f496a7068


[ROCm/clr commit: a8c2375fe1]
This commit is contained in:
Satyanvesh Dittakavi
2023-02-28 16:22:15 +00:00
parent 2bcde46b12
commit 8947306ffd
4 changed files with 79 additions and 73 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h
+68 -68
View File
@@ -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 <hip/amd_detail/amd_hip_common.h>
#include "hip/amd_detail/host_defines.h"
#if defined(__HIPCC_RTC__)
#define __HOST_DEVICE__ __device__
#else
#define __HOST_DEVICE__ __host__ __device__
#include <hip/amd_detail/amd_hip_common.h>
#include "hip/amd_detail/host_defines.h"
#include <assert.h>
#if defined(__cplusplus)
#include <algorithm>
@@ -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
projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h
+4 -2
View File
@@ -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
projects/clr/hipamd/src/hiprtc/CMakeLists.txt
+3 -1
View File
@@ -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 $<TARGET_FILE:clang> -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}
projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake
+4 -2
View File
@@ -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\")")