Started adding native half math library support
1. Removed HIP_EXPERIMENTAL env variable so that device code will be accessed from LLVM IR
2. Removed soft support from headers and moved to hip_fp16.cpp
3. Added LLVM IR + inline asm to hip_ir.ll
4. Added test for fp16
5. Added barriers for hcc 3.5 and hcc 4.0 for half support
a. Which means, hcc 4.0 can parse __fp16 but hcc 3.5 cant
b. HCC 4.0 code is implemented now, hcc 3.5 will be added later
Change-Id: Ic37859b2688ebb02e168bab643d1882bf4727952
[ROCm/hip commit: d180fdaae0]
Este cometimento está contido em:
Aditya Atluri
@@ -25,213 +25,81 @@ THE SOFTWARE.
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#if 0
|
||||
#if __clang_major__ == 4
|
||||
|
||||
typedef __fp16 __half;
|
||||
|
||||
typedef struct __attribute__((aligned(4))){
|
||||
int a;
|
||||
union {
|
||||
__half p[2];
|
||||
unsigned int q;
|
||||
};
|
||||
} __half2;
|
||||
|
||||
extern "C" __half __hip_hadd_gfx803(__half a, __half b);
|
||||
extern "C" __half __hip_hfma_gfx803(__half a, __half b);
|
||||
extern "C" __half __hip_hmul_gfx803(__half a, __half b);
|
||||
extern "C" __half __hip_hsub_gfx803(__half a, __half b);
|
||||
extern "C" __half __hip_hc_ir_hadd_half(__half, __half);
|
||||
extern "C" __half __hip_hc_ir_hfma_half(__half, __half, __half);
|
||||
extern "C" __half __hip_hc_ir_hmul_half(__half, __half);
|
||||
extern "C" __half __hip_hc_ir_hsub_half(__half, __half);
|
||||
|
||||
extern "C" int __hip_hadd2_gfx803(int a, int b);
|
||||
extern "C" int __hip_hfma2_gfx803(int a, int b);
|
||||
extern "C" int __hip_hmul2_gfx803(int a, int b);
|
||||
extern "C" int __hip_hsub2_gfx803(int a, int b);
|
||||
|
||||
__device__ inline __half __hadd(__half a, __half b) {
|
||||
return __hip_hadd_gfx803(a, b);
|
||||
__device__ static inline __half __hadd(const __half a, const __half b) {
|
||||
return __hip_hc_ir_hadd_half(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hadd_sat(__half a, __half b) {
|
||||
return __hip_hadd_gfx803(a, b);
|
||||
__device__ static inline __half __hadd_sat(__half a, __half b) {
|
||||
return __hip_hc_ir_hadd_half(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hfma(__half a, __half b) {
|
||||
return __hip_hfma_gfx803(a, b);
|
||||
__device__ static inline __half __hfma(__half a, __half b, __half c) {
|
||||
return __hip_hc_ir_hfma_half(a, b, c);
|
||||
}
|
||||
|
||||
__device__ inline __half __hfma_sat(__half a, __half b) {
|
||||
return __hip_hfma_gfx803(a, b);
|
||||
__device__ static inline __half __hfma_sat(__half a, __half b, __half c) {
|
||||
return __hip_hc_ir_hfma_half(a, b, c);
|
||||
}
|
||||
|
||||
__device__ inline __half __hmul(__half a, __half b) {
|
||||
return __hip_hmul_gfx803(a, b);
|
||||
__device__ static inline __half __hmul(__half a, __half b) {
|
||||
return __hip_hc_ir_hmul_half(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hmul_sat(__half a, __half b) {
|
||||
return __hip_hmul_gfx803(a, b);
|
||||
__device__ static inline __half __hmul_sat(__half a, __half b) {
|
||||
return __hip_hc_ir_hmul_half(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hsub(__half a, __half b) {
|
||||
return __hip_hsub_gfx803(a, b);
|
||||
__device__ static inline __half __hneg(__half a) {
|
||||
return -a;
|
||||
}
|
||||
|
||||
__device__ inline __half __hsub_sat(__half a, __half b) {
|
||||
return __hip_hsub_gfx803(a, b);
|
||||
__device__ static inline __half __hsub(__half a, __half b) {
|
||||
return __hip_hc_ir_hsub_half(a, b);
|
||||
}
|
||||
|
||||
|
||||
__device__ inline __half2 __hadd2(__half2 a, __half2 b) {
|
||||
__half2 ret;
|
||||
ret.a = __hip_hadd2_gfx803(a.a, b.a);
|
||||
return ret;
|
||||
__device__ static inline __half __hsub_sat(__half a, __half b) {
|
||||
return __hip_hc_ir_hsub_half(a, b);
|
||||
}
|
||||
|
||||
#else
|
||||
__device__ static inline __half hdiv(__half a, __half b) {
|
||||
return a/b;
|
||||
}
|
||||
|
||||
typedef struct{
|
||||
#endif
|
||||
|
||||
#if __clang_major__ == 3
|
||||
|
||||
typedef struct {
|
||||
unsigned x: 16;
|
||||
} __half;
|
||||
|
||||
|
||||
typedef struct __attribute__((aligned(4))){
|
||||
__half p,q;
|
||||
union {
|
||||
__half p[2];
|
||||
unsigned int q;
|
||||
};
|
||||
} __half2;
|
||||
|
||||
typedef __half half;
|
||||
typedef __half2 half2;
|
||||
|
||||
/*
|
||||
Arithmetic functions
|
||||
*/
|
||||
|
||||
__device__ __half __hadd(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hadd_sat(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hfma(const __half a, const __half b, const __half c);
|
||||
|
||||
__device__ __half __hfma_sat(const __half a, const __half b, const __half c);
|
||||
|
||||
__device__ __half __hmul(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hmul_sat(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hneq(const __half a);
|
||||
|
||||
__device__ __half __hsub(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hsub_sat(const __half a, const __half b);
|
||||
|
||||
|
||||
|
||||
/*
|
||||
Half2 Arithmetic Instructions
|
||||
*/
|
||||
|
||||
__device__ __half2 __hadd2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hadd2_sat(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hfma2(const __half2 a, const __half2 b, const __half2 c);
|
||||
|
||||
__device__ __half2 __hfma2_sat(const __half2 a, const __half2 b, const __half2 c);
|
||||
|
||||
__device__ __half2 __hmul2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hmul2_sat(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hneq2(const __half2 a);
|
||||
|
||||
__device__ __half2 __hsub2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hsub2_sat(const __half2 a, const __half2 b);
|
||||
|
||||
/*
|
||||
Half Cmps
|
||||
*/
|
||||
|
||||
__device__ bool __heq(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hge(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hgt(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hisinf(const __half a);
|
||||
|
||||
__device__ bool __hisnan(const __half a);
|
||||
|
||||
__device__ bool __hle(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hlt(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hne(const __half a, const __half b);
|
||||
|
||||
/*
|
||||
Half2 Cmps
|
||||
*/
|
||||
|
||||
__device__ bool __hbeq2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hbge2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hbgt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hble2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hblt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hbne2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __heq2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hge2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hgt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hisnan2(const __half2 a);
|
||||
|
||||
__device__ __half2 __hle2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hlt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hne2(const __half2 a, const __half2 b);
|
||||
|
||||
|
||||
/*
|
||||
Half Cnvs and Data Mvmnt
|
||||
*/
|
||||
|
||||
__device__ __half2 __float22half2_rn(const float2 a);
|
||||
|
||||
__device__ __half __float2half(const float a);
|
||||
|
||||
__device__ __half2 __float2half2_rn(const float a);
|
||||
|
||||
__device__ __half2 __floats2half2_rn(const float a, const float b);
|
||||
|
||||
__device__ float2 __half22float2(const __half2 a);
|
||||
|
||||
__device__ float __half2float(const __half a);
|
||||
|
||||
__device__ __half2 __half2half2(const __half a);
|
||||
|
||||
__device__ __half2 __halves2half2(const __half a, const __half b);
|
||||
|
||||
__device__ float __high2float(const __half2 a);
|
||||
|
||||
__device__ __half __high2half(const __half2 a);
|
||||
|
||||
__device__ __half2 __high2half2(const __half2 a);
|
||||
|
||||
__device__ __half2 __highs2half2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ float __low2float(const __half2 a);
|
||||
|
||||
__device__ __half __low2half(const __half2 a);
|
||||
|
||||
__device__ __half2 __low2half2(const __half2 a);
|
||||
|
||||
__device__ __half2 __lows2half2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __lowhigh2highlow(const __half2 a);
|
||||
|
||||
__device__ __half2 __low2half2(const __half2 a, const __half2 b);
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
#endif
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador