Aditya Atluri
01b66dd998
1. Split hip_ir.ll to hip_hc.ll and hip_hc_gfx803.ll a. hip_hc.ll contains arch generic ir implementations b. hip_hc_gfx803.ll contains gfx803 (fiji, polaris) specific ir 2. HIPCC can now parse --amdgpu-target=*. a. Usage: hipcc --amdgpu-target=gfx803 --amdgpu-target=gfx701 b. TODO: Convert to --amdgpu-target=gfx803,gfx701 3. With LLC in HCC able to generate native f16 isa, removed inline half asm math ops 4. Fixed threadfence and threadfence_block to use functions in rocdl Change-Id: Ic9a9e3e04139b0d75d2c2a263c030ca77adc1019
487 wiersze
15 KiB
C
487 wiersze
15 KiB
C
/*
|
|
Copyright (c) 2015-2017 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.
|
|
*/
|
|
|
|
#ifndef HIP_HCC_DETAIL_FP16_H
|
|
#define HIP_HCC_DETAIL_FP16_H
|
|
|
|
#include "hip/hip_runtime.h"
|
|
|
|
#if __clang_major__ == 4
|
|
|
|
typedef __fp16 __half;
|
|
|
|
typedef struct __attribute__((aligned(4))){
|
|
union {
|
|
__half p[2];
|
|
unsigned int q;
|
|
};
|
|
} __half2;
|
|
|
|
/*
|
|
Half Arithmetic Functions
|
|
*/
|
|
__device__ __half __hadd(const __half a, const __half b);
|
|
__device__ __half __hadd_sat(__half a, __half b);
|
|
__device__ __half __hfma(__half a, __half b, __half c);
|
|
__device__ __half __hfma_sat(__half a, __half b, __half c);
|
|
__device__ __half __hmul(__half a, __half b);
|
|
__device__ __half __hmul_sat(__half a, __half b);
|
|
__device__ __half __hneg(__half a);
|
|
__device__ __half __hsub(__half a, __half b);
|
|
__device__ __half __hsub_sat(__half a, __half b);
|
|
__device__ __half hdiv(__half a, __half b);
|
|
|
|
/*
|
|
Half2 Arithmetic Functions
|
|
*/
|
|
|
|
__device__ static __half2 __hadd2(__half2 a, __half2 b);
|
|
__device__ static __half2 __hadd2_sat(__half2 a, __half2 b);
|
|
__device__ static __half2 __hfma2(__half2 a, __half2 b, __half2 c);
|
|
__device__ static __half2 __hfma2_sat(__half2 a, __half2 b, __half2 c);
|
|
__device__ static __half2 __hmul2(__half2 a, __half2 b);
|
|
__device__ static __half2 __hmul2_sat(__half2 a, __half2 b);
|
|
__device__ static __half2 __hsub2(__half2 a, __half2 b);
|
|
__device__ static __half2 __hneg2(__half2 a);
|
|
__device__ static __half2 __hsub2_sat(__half2 a, __half2 b);
|
|
__device__ static __half2 h2div(__half2 a, __half2 b);
|
|
|
|
/*
|
|
Half Comparision Functions
|
|
*/
|
|
|
|
__device__ bool __heq(__half a, __half b);
|
|
__device__ bool __hge(__half a, __half b);
|
|
__device__ bool __hgt(__half a, __half b);
|
|
__device__ bool __hisinf(__half a);
|
|
__device__ bool __hisnan(__half a);
|
|
__device__ bool __hle(__half a, __half b);
|
|
__device__ bool __hlt(__half a, __half b);
|
|
__device__ bool __hne(__half a, __half b);
|
|
|
|
/*
|
|
Half2 Comparision Functions
|
|
*/
|
|
|
|
__device__ bool __hbeq2(__half2 a, __half2 b);
|
|
__device__ bool __hbge2(__half2 a, __half2 b);
|
|
__device__ bool __hbgt2(__half2 a, __half2 b);
|
|
__device__ bool __hble2(__half2 a, __half2 b);
|
|
__device__ bool __hblt2(__half2 a, __half2 b);
|
|
__device__ bool __hbne2(__half2 a, __half2 b);
|
|
__device__ __half2 __heq2(__half2 a, __half2 b);
|
|
__device__ __half2 __hge2(__half2 a, __half2 b);
|
|
__device__ __half2 __hgt2(__half2 a, __half2 b);
|
|
__device__ __half2 __hisnan2(__half2 a);
|
|
__device__ __half2 __hle2(__half2 a, __half2 b);
|
|
__device__ __half2 __hlt2(__half2 a, __half2 b);
|
|
__device__ __half2 __hne2(__half2 a, __half2 b);
|
|
|
|
/*
|
|
Half Math Functions
|
|
*/
|
|
|
|
__device__ static __half hceil(const __half h);
|
|
__device__ static __half hcos(const __half h);
|
|
__device__ static __half hexp(const __half h);
|
|
__device__ static __half hexp10(const __half h);
|
|
__device__ static __half hexp2(const __half h);
|
|
__device__ static __half hfloor(const __half h);
|
|
__device__ static __half hlog(const __half h);
|
|
__device__ static __half hlog10(const __half h);
|
|
__device__ static __half hlog2(const __half h);
|
|
//__device__ static __half hrcp(const __half h);
|
|
__device__ static __half hrint(const __half h);
|
|
__device__ static __half hsin(const __half h);
|
|
__device__ static __half hsqrt(const __half a);
|
|
__device__ static __half htrunc(const __half a);
|
|
|
|
/*
|
|
Half2 Math Functions
|
|
*/
|
|
|
|
__device__ static __half2 h2ceil(const __half2 h);
|
|
__device__ static __half2 h2exp(const __half2 h);
|
|
__device__ static __half2 h2exp10(const __half2 h);
|
|
__device__ static __half2 h2exp2(const __half2 h);
|
|
__device__ static __half2 h2floor(const __half2 h);
|
|
__device__ static __half2 h2log(const __half2 h);
|
|
__device__ static __half2 h2log10(const __half2 h);
|
|
__device__ static __half2 h2log2(const __half2 h);
|
|
__device__ static __half2 h2rcp(const __half2 h);
|
|
__device__ static __half2 h2rsqrt(const __half2 h);
|
|
__device__ static __half2 h2sin(const __half2 h);
|
|
__device__ static __half2 h2sqrt(const __half2 h);
|
|
|
|
/*
|
|
Half Conversion And Data Movement
|
|
*/
|
|
|
|
__device__ __half2 __float22half2_rn(const float2 a);
|
|
__device__ __half __float2half(const float a);
|
|
__device__ __half2 __float2half2_rn(const float a);
|
|
__device__ __half __float2half_rd(const float a);
|
|
__device__ __half __float2half_rn(const float a);
|
|
__device__ __half __float2half_ru(const float a);
|
|
__device__ __half __float2half_rz(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__ int __half2int_rd(__half h);
|
|
__device__ int __half2int_rn(__half h);
|
|
__device__ int __half2int_ru(__half h);
|
|
__device__ int __half2int_rz(__half h);
|
|
__device__ long long int __half2ll_rd(__half h);
|
|
__device__ long long int __half2ll_rn(__half h);
|
|
__device__ long long int __half2ll_ru(__half h);
|
|
__device__ long long int __half2ll_rz(__half h);
|
|
__device__ short __half2short_rd(__half h);
|
|
__device__ short __half2short_rn(__half h);
|
|
__device__ short __half2short_ru(__half h);
|
|
__device__ short __half2short_rz(__half h);
|
|
__device__ unsigned int __half2uint_rd(__half h);
|
|
__device__ unsigned int __half2uint_rn(__half h);
|
|
__device__ unsigned int __half2uint_ru(__half h);
|
|
__device__ unsigned int __half2uint_rz(__half h);
|
|
__device__ unsigned long long int __half2ull_rd(__half h);
|
|
__device__ unsigned long long int __half2ull_rn(__half h);
|
|
__device__ unsigned long long int __half2ull_ru(__half h);
|
|
__device__ unsigned long long int __half2ull_rz(__half h);
|
|
__device__ unsigned short int __half2ushort_rd(__half h);
|
|
__device__ unsigned short int __half2ushort_rn(__half h);
|
|
__device__ unsigned short int __half2ushort_ru(__half h);
|
|
__device__ unsigned short int __half2ushort_rz(__half h);
|
|
__device__ short int __half_as_short(const __half h);
|
|
__device__ unsigned short int __half_as_ushort(const __half h);
|
|
__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__ __half __int2half_rd(int i);
|
|
__device__ __half __int2half_rn(int i);
|
|
__device__ __half __int2half_ru(int i);
|
|
__device__ __half __int2half_rz(int i);
|
|
__device__ __half __ll2half_rd(long long int i);
|
|
__device__ __half __ll2half_rn(long long int i);
|
|
__device__ __half __ll2half_ru(long long int i);
|
|
__device__ __half __ll2half_rz(long long int i);
|
|
__device__ float __low2float(const __half2 a);
|
|
|
|
__device__ __half __low2half(const __half2 a);
|
|
__device__ __half2 __low2half2(const __half2 a, const __half2 b);
|
|
__device__ __half2 __low2half2(const __half2 a);
|
|
__device__ __half2 __lowhigh2highlow(const __half2 a);
|
|
__device__ __half2 __lows2half2(const __half2 a, const __half2 b);
|
|
__device__ __half __short2half_rd(short int i);
|
|
__device__ __half __short2half_rn(short int i);
|
|
__device__ __half __short2half_ru(short int i);
|
|
__device__ __half __short2half_rz(short int i);
|
|
__device__ __half __uint2half_rd(unsigned int i);
|
|
__device__ __half __uint2half_rn(unsigned int i);
|
|
__device__ __half __uint2half_ru(unsigned int i);
|
|
__device__ __half __uint2half_rz(unsigned int i);
|
|
__device__ __half __ull2half_rd(unsigned long long int i);
|
|
__device__ __half __ull2half_rn(unsigned long long int i);
|
|
__device__ __half __ull2half_ru(unsigned long long int i);
|
|
__device__ __half __ull2half_rz(unsigned long long int i);
|
|
__device__ __half __ushort2half_rd(unsigned short int i);
|
|
__device__ __half __ushort2half_rn(unsigned short int i);
|
|
__device__ __half __ushort2half_ru(unsigned short int i);
|
|
__device__ __half __ushort2half_rz(unsigned short int i);
|
|
__device__ __half __ushort_as_half(const unsigned short int i);
|
|
|
|
extern "C" int __hip_hc_ir_hadd2_int(int, int);
|
|
extern "C" int __hip_hc_ir_hfma2_int(int, int, int);
|
|
extern "C" int __hip_hc_ir_hmul2_int(int, int);
|
|
extern "C" int __hip_hc_ir_hsub2_int(int, int);
|
|
|
|
extern "C" __half __hip_hc_ir_hceil_half(__half) __asm("llvm.ceil.f16");
|
|
extern "C" __half __hip_hc_ir_hcos_half(__half) __asm("llvm.cos.f16");
|
|
extern "C" __half __hip_hc_ir_hexp2_half(__half) __asm("llvm.exp2.f16");
|
|
extern "C" __half __hip_hc_ir_hfloor_half(__half) __asm("llvm.floor.f16");
|
|
extern "C" __half __hip_hc_ir_hlog2_half(__half) __asm("llvm.log2.f16");
|
|
extern "C" __half __hip_hc_ir_hrcp_half(__half) __asm("llvm.amdgcn.rcp.f16");
|
|
extern "C" __half __hip_hc_ir_hrint_half(__half) __asm("llvm.rint.f16");
|
|
extern "C" __half __hip_hc_ir_hrsqrt_half(__half) __asm("llvm.sqrt.f16");
|
|
extern "C" __half __hip_hc_ir_hsin_half(__half) __asm("llvm.sin.f16");
|
|
extern "C" __half __hip_hc_ir_hsqrt_half(__half) __asm("llvm.sqrt.f16");
|
|
extern "C" __half __hip_hc_ir_htrunc_half(__half) __asm("llvm.trunc.f16");
|
|
|
|
extern "C" int __hip_hc_ir_h2ceil_int(int);
|
|
extern "C" int __hip_hc_ir_h2cos_int(int);
|
|
extern "C" int __hip_hc_ir_h2exp2_int(int);
|
|
extern "C" int __hip_hc_ir_h2floor_int(int);
|
|
extern "C" int __hip_hc_ir_h2log2_int(int);
|
|
extern "C" int __hip_hc_ir_h2rcp_int(int);
|
|
extern "C" int __hip_hc_ir_h2rsqrt_int(int);
|
|
extern "C" int __hip_hc_ir_h2sin_int(int);
|
|
extern "C" int __hip_hc_ir_h2sqrt_int(int);
|
|
extern "C" int __hip_hc_ir_h2trunc_int(int);
|
|
|
|
/*
|
|
Half2 Arithmetic Functions
|
|
*/
|
|
|
|
__device__ static inline __half2 __hadd2(__half2 a, __half2 b) {
|
|
__half2 c;
|
|
c.q = __hip_hc_ir_hadd2_int(a.q, b.q);
|
|
return c;
|
|
}
|
|
|
|
__device__ static inline __half2 __hadd2_sat(__half2 a, __half2 b) {
|
|
__half2 c;
|
|
c.q = __hip_hc_ir_hadd2_int(a.q, b.q);
|
|
return c;
|
|
}
|
|
|
|
__device__ static inline __half2 __hfma2(__half2 a, __half2 b, __half2 c) {
|
|
__half2 d;
|
|
d.q = __hip_hc_ir_hfma2_int(a.q, b.q, c.q);
|
|
return d;
|
|
}
|
|
|
|
__device__ static inline __half2 __hfma2_sat(__half2 a, __half2 b, __half2 c) {
|
|
__half2 d;
|
|
d.q = __hip_hc_ir_hfma2_int(a.q, b.q, c.q);
|
|
return d;
|
|
}
|
|
|
|
__device__ static inline __half2 __hmul2(__half2 a, __half2 b) {
|
|
__half2 c;
|
|
c.q = __hip_hc_ir_hmul2_int(a.q, b.q);
|
|
return c;
|
|
}
|
|
|
|
__device__ static inline __half2 __hmul2_sat(__half2 a, __half2 b) {
|
|
__half2 c;
|
|
c.q = __hip_hc_ir_hmul2_int(a.q, b.q);
|
|
return c;
|
|
}
|
|
|
|
__device__ static inline __half2 __hsub2(__half2 a, __half2 b) {
|
|
__half2 c;
|
|
c.q = __hip_hc_ir_hsub2_int(a.q, b.q);
|
|
return c;
|
|
}
|
|
|
|
__device__ static inline __half2 __hneg2(__half2 a) {
|
|
__half2 c;
|
|
c.p[0] = - a.p[0];
|
|
c.p[1] = - a.p[1];
|
|
return c;
|
|
}
|
|
|
|
__device__ static inline __half2 __hsub2_sat(__half2 a, __half2 b) {
|
|
__half2 c;
|
|
c.q = __hip_hc_ir_hsub2_int(a.q, b.q);
|
|
return c;
|
|
}
|
|
|
|
__device__ static inline __half2 h2div(__half2 a, __half2 b) {
|
|
__half2 c;
|
|
c.p[0] = a.p[0] / b.p[0];
|
|
c.p[1] = a.p[1] / b.p[1];
|
|
return c;
|
|
}
|
|
|
|
|
|
__device__ static inline __half hceil(const __half h) {
|
|
return __hip_hc_ir_hceil_half(h);
|
|
}
|
|
|
|
__device__ static inline __half hcos(const __half h) {
|
|
return __hip_hc_ir_hcos_half(h);
|
|
}
|
|
|
|
__device__ static inline __half hexp(const __half h) {
|
|
return __hip_hc_ir_hexp2_half(__hmul(h, 1.442694));
|
|
}
|
|
|
|
__device__ static inline __half hexp10(const __half h) {
|
|
return __hip_hc_ir_hexp2_half(__hmul(h, 3.3219281));
|
|
}
|
|
|
|
__device__ static inline __half hexp2(const __half h) {
|
|
return __hip_hc_ir_hexp2_half(h);
|
|
}
|
|
|
|
__device__ static inline __half hfloor(const __half h) {
|
|
return __hip_hc_ir_hfloor_half(h);
|
|
}
|
|
|
|
__device__ static inline __half hlog(const __half h) {
|
|
return __hmul(__hip_hc_ir_hlog2_half(h), 0.693147);
|
|
}
|
|
|
|
__device__ static inline __half hlog10(const __half h) {
|
|
return __hmul(__hip_hc_ir_hlog2_half(h), 0.301029);
|
|
}
|
|
|
|
__device__ static inline __half hlog2(const __half h) {
|
|
return __hip_hc_ir_hlog2_half(h);
|
|
}
|
|
/*
|
|
__device__ static inline __half hrcp(const __half h) {
|
|
return __hip_hc_ir_hrcp_half(h);
|
|
}
|
|
*/
|
|
__device__ static inline __half hrint(const __half h) {
|
|
return __hip_hc_ir_hrint_half(h);
|
|
}
|
|
|
|
__device__ static inline __half hrsqrt(const __half h) {
|
|
return __hip_hc_ir_hrsqrt_half(h);
|
|
}
|
|
|
|
__device__ static inline __half hsin(const __half h) {
|
|
return __hip_hc_ir_hsin_half(h);
|
|
}
|
|
|
|
__device__ static inline __half hsqrt(const __half a) {
|
|
return __hip_hc_ir_hsqrt_half(a);
|
|
}
|
|
|
|
__device__ static inline __half htrunc(const __half a) {
|
|
return __hip_hc_ir_htrunc_half(a);
|
|
}
|
|
|
|
/*
|
|
Half2 Math Operations
|
|
*/
|
|
|
|
__device__ static inline __half2 h2ceil(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2ceil_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2cos(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2cos_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2exp(const __half2 h) {
|
|
__half2 factor;
|
|
factor.p[0] = 1.442694;
|
|
factor.p[1] = 1.442694;
|
|
factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q));
|
|
return factor;
|
|
}
|
|
|
|
__device__ static inline __half2 h2exp10(const __half2 h) {
|
|
__half2 factor;
|
|
factor.p[0] = 3.3219281;
|
|
factor.p[1] = 3.3219281;
|
|
factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q));
|
|
return factor;
|
|
}
|
|
|
|
__device__ static inline __half2 h2exp2(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2exp2_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2floor(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2floor_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2log(const __half2 h) {
|
|
__half2 factor;
|
|
factor.p[0] = 0.693147;
|
|
factor.p[1] = 0.693147;
|
|
factor. q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q);
|
|
return factor;
|
|
}
|
|
|
|
__device__ static inline __half2 h2log10(const __half2 h) {
|
|
__half2 factor;
|
|
factor.p[0] = 0.301029;
|
|
factor.p[1] = 0.301029;
|
|
factor.q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q);
|
|
return factor;
|
|
}
|
|
__device__ static inline __half2 h2log2(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2log2_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2rcp(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2rcp_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2rsqrt(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2rsqrt_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2sin(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2sin_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2sqrt(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2sqrt_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
__device__ static inline __half2 h2trunc(const __half2 h) {
|
|
__half2 a;
|
|
a.q = __hip_hc_ir_h2trunc_int(h.q);
|
|
return a;
|
|
}
|
|
|
|
|
|
|
|
|
|
#endif
|
|
|
|
#if __clang_major__ == 3
|
|
|
|
typedef struct {
|
|
unsigned x: 16;
|
|
} __half;
|
|
|
|
typedef struct __attribute__((aligned(4))){
|
|
union {
|
|
__half p[2];
|
|
unsigned int q;
|
|
};
|
|
} __half2;
|
|
|
|
|
|
#endif
|
|
|
|
|
|
#endif
|