changed vector types to make sure it generate proper llvm vector types
Change-Id: I6c4616dae137dc4eac35e5827dc5b7f3251e0247
Этот коммит содержится в:
Aditya Atluri
@@ -25,17 +25,6 @@ THE SOFTWARE.
|
||||
|
||||
#include "hip/hcc_detail/hip_vector_types.h"
|
||||
|
||||
#if __clang_major__ > 3
|
||||
|
||||
typedef __fp16 __half;
|
||||
|
||||
typedef struct __attribute__((aligned(4))){
|
||||
union {
|
||||
__half p[2];
|
||||
unsigned int q;
|
||||
};
|
||||
} __half2;
|
||||
|
||||
typedef __half half;
|
||||
typedef __half2 half2;
|
||||
|
||||
@@ -214,10 +203,10 @@ __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" __half2 __hip_hc_ir_hadd2_int(__half2, __half2);
|
||||
extern "C" __half2 __hip_hc_ir_hfma2_int(__half2, __half2, __half2);
|
||||
extern "C" __half2 __hip_hc_ir_hmul2_int(__half2, __half2);
|
||||
extern "C" __half2 __hip_hc_ir_hsub2_int(__half2, __half2);
|
||||
|
||||
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");
|
||||
@@ -231,16 +220,16 @@ 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);
|
||||
extern "C" __half2 __hip_hc_ir_h2ceil_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2cos_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2exp2_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2floor_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2log2_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2rcp_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2rsqrt_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2sin_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2sqrt_int(__half2);
|
||||
extern "C" __half2 __hip_hc_ir_h2trunc_int(__half2);
|
||||
|
||||
/*
|
||||
Half2 Arithmetic Functions
|
||||
@@ -248,63 +237,63 @@ extern "C" int __hip_hc_ir_h2trunc_int(int);
|
||||
|
||||
__device__ static inline __half2 __hadd2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.q = __hip_hc_ir_hadd2_int(a.q, b.q);
|
||||
c.xy = __hip_hc_ir_hadd2_int(a.xy, b.xy);
|
||||
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);
|
||||
c.xy = __hip_hc_ir_hadd2_int(a.xy, b.xy);
|
||||
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);
|
||||
d.xy = __hip_hc_ir_hfma2_int(a.xy, b.xy, c.xy);
|
||||
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);
|
||||
d.xy = __hip_hc_ir_hfma2_int(a.xy, b.xy, c.xy);
|
||||
return d;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 __hmul2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.q = __hip_hc_ir_hmul2_int(a.q, b.q);
|
||||
c.xy = __hip_hc_ir_hmul2_int(a.xy, b.xy);
|
||||
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);
|
||||
c.xy = __hip_hc_ir_hmul2_int(a.xy, b.xy);
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 __hsub2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.q = __hip_hc_ir_hsub2_int(a.q, b.q);
|
||||
c.xy = __hip_hc_ir_hsub2_int(a.xy, b.xy);
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 __hneg2(__half2 a) {
|
||||
__half2 c;
|
||||
c.p[0] = - a.p[0];
|
||||
c.p[1] = - a.p[1];
|
||||
c.x = - a.x;
|
||||
c.y = - a.y;
|
||||
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);
|
||||
c.xy = __hip_hc_ir_hsub2_int(a.xy, b.xy);
|
||||
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];
|
||||
c.x = a.x / b.x;
|
||||
c.y = a.y / b.y;
|
||||
return c;
|
||||
}
|
||||
|
||||
@@ -375,112 +364,94 @@ Half2 Math Operations
|
||||
|
||||
__device__ static inline __half2 h2ceil(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2ceil_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2ceil_int(h.xy);
|
||||
return a;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2cos(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2cos_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2cos_int(h.xy);
|
||||
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));
|
||||
factor.x = 1.442694;
|
||||
factor.y = 1.442694;
|
||||
factor.xy = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.xy, factor.xy));
|
||||
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));
|
||||
factor.x = 3.3219281;
|
||||
factor.y = 3.3219281;
|
||||
factor.xy = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.xy, factor.xy));
|
||||
return factor;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2exp2(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2exp2_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2exp2_int(h.xy);
|
||||
return a;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2floor(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2floor_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2floor_int(h.xy);
|
||||
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);
|
||||
factor.x = 0.693147;
|
||||
factor.y = 0.693147;
|
||||
factor.xy = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.xy), factor.xy);
|
||||
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);
|
||||
factor.x = 0.301029;
|
||||
factor.y = 0.301029;
|
||||
factor.xy = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.xy), factor.xy);
|
||||
return factor;
|
||||
}
|
||||
__device__ static inline __half2 h2log2(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2log2_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2log2_int(h.xy);
|
||||
return a;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2rcp(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2rcp_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2rcp_int(h.xy);
|
||||
return a;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2rsqrt(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2rsqrt_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2rsqrt_int(h.xy);
|
||||
return a;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2sin(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2sin_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2sin_int(h.xy);
|
||||
return a;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2sqrt(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2sqrt_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2sqrt_int(h.xy);
|
||||
return a;
|
||||
}
|
||||
|
||||
__device__ static inline __half2 h2trunc(const __half2 h) {
|
||||
__half2 a;
|
||||
a.q = __hip_hc_ir_h2trunc_int(h.q);
|
||||
a.xy = __hip_hc_ir_h2trunc_int(h.xy);
|
||||
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
|
||||
|
||||
Разница между файлами не показана из-за своего большого размера
Загрузить разницу
+50
-392
@@ -90,11 +90,11 @@ __device__ bool __hgt(__half a, __half b) {
|
||||
}
|
||||
|
||||
__device__ bool __hisinf(__half a) {
|
||||
return a == __hInfValue.h ? true : false;
|
||||
return a == HINF ? true : false;
|
||||
}
|
||||
|
||||
__device__ bool __hisnan(__half a) {
|
||||
return a > __hInfValue.h ? true : false;
|
||||
return a > HINF ? true : false;
|
||||
}
|
||||
|
||||
__device__ bool __hle(__half a, __half b) {
|
||||
@@ -114,75 +114,75 @@ Half2 Comparision Functions
|
||||
*/
|
||||
|
||||
__device__ bool __hbeq2(__half2 a, __half2 b) {
|
||||
return (a.p[0] == b.p[0] ? true : false) && (a.p[1] == b.p[1] ? true : false);
|
||||
return (a.x == b.x ? true : false) && (a.y == b.y ? true : false);
|
||||
}
|
||||
|
||||
__device__ bool __hbge2(__half2 a, __half2 b) {
|
||||
return (a.p[0] >= b.p[0] ? true : false) && (a.p[1] >= b.p[1] ? true : false);
|
||||
return (a.x >= b.x ? true : false) && (a.y >= b.y ? true : false);
|
||||
}
|
||||
|
||||
__device__ bool __hbgt2(__half2 a, __half2 b) {
|
||||
return (a.p[0] > b.p[0] ? true : false) && (a.p[1] > b.p[1] ? true : false);
|
||||
return (a.x > b.x ? true : false) && (a.y > b.y ? true : false);
|
||||
}
|
||||
|
||||
__device__ bool __hble2(__half2 a, __half2 b) {
|
||||
return (a.p[0] <= b.p[0] ? true : false) && (a.p[1] <= b.p[1] ? true : false);
|
||||
return (a.x <= b.x ? true : false) && (a.y <= b.y ? true : false);
|
||||
}
|
||||
|
||||
__device__ bool __hblt2(__half2 a, __half2 b) {
|
||||
return (a.p[0] < b.p[0] ? true : false) && (a.p[1] < b.p[1] ? true : false);
|
||||
return (a.x < b.x ? true : false) && (a.y < b.y ? true : false);
|
||||
}
|
||||
|
||||
__device__ bool __hbne2(__half2 a, __half2 b) {
|
||||
return (a.p[0] != b.p[0] ? true : false) && (a.p[1] != b.p[1] ? true : false);
|
||||
return (a.x != b.x ? true : false) && (a.y != b.y ? true : false);
|
||||
}
|
||||
|
||||
__device__ __half2 __heq2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = (a.p[0] == b.p[0]) ? (__half)1 : (__half)0;
|
||||
c.p[1] = (a.p[1] == b.p[1]) ? (__half)1 : (__half)0;
|
||||
c.x = (a.x == b.x) ? (__half)1 : (__half)0;
|
||||
c.y = (a.y == b.y) ? (__half)1 : (__half)0;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ __half2 __hge2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = (a.p[0] >= b.p[0]) ? (__half)1 : (__half)0;
|
||||
c.p[1] = (a.p[1] >= b.p[1]) ? (__half)1 : (__half)0;
|
||||
c.x = (a.x >= b.x) ? (__half)1 : (__half)0;
|
||||
c.y = (a.y >= b.y) ? (__half)1 : (__half)0;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ __half2 __hgt2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = (a.p[0] > b.p[0]) ? (__half)1 : (__half)0;
|
||||
c.p[1] = (a.p[1] > b.p[1]) ? (__half)1 : (__half)0;
|
||||
c.x = (a.x > b.x) ? (__half)1 : (__half)0;
|
||||
c.y = (a.y > b.y) ? (__half)1 : (__half)0;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ __half2 __hisnan2(__half2 a) {
|
||||
__half2 c;
|
||||
c.p[0] = (a.p[0] > __hInfValue.h) ? (__half)1 : (__half)0;
|
||||
c.p[1] = (a.p[1] > __hInfValue.h) ? (__half)1 : (__half)0;
|
||||
c.x = (a.x > HINF) ? (__half)1 : (__half)0;
|
||||
c.y = (a.y > HINF) ? (__half)1 : (__half)0;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ __half2 __hle2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = (a.p[0] <= b.p[0]) ? (__half)1 : (__half)0;
|
||||
c.p[1] = (a.p[1] <= b.p[1]) ? (__half)1 : (__half)0;
|
||||
c.x = (a.x <= b.x) ? (__half)1 : (__half)0;
|
||||
c.y = (a.y <= b.y) ? (__half)1 : (__half)0;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ __half2 __hlt2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = (a.p[0] < b.p[0]) ? (__half)1 : (__half)0;
|
||||
c.p[1] = (a.p[1] < b.p[1]) ? (__half)1 : (__half)0;
|
||||
c.x = (a.x < b.x) ? (__half)1 : (__half)0;
|
||||
c.y = (a.y < b.y) ? (__half)1 : (__half)0;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ __half2 __hne2(__half2 a, __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = (a.p[0] != b.p[0]) ? (__half)1 : (__half)0;
|
||||
c.p[1] = (a.p[1] != b.p[1]) ? (__half)1 : (__half)0;
|
||||
c.x = (a.x != b.x) ? (__half)1 : (__half)0;
|
||||
c.y = (a.y != b.y) ? (__half)1 : (__half)0;
|
||||
return c;
|
||||
}
|
||||
|
||||
@@ -191,8 +191,8 @@ Conversion instructions
|
||||
*/
|
||||
__device__ __half2 __float22half2_rn(const float2 a) {
|
||||
__half2 b;
|
||||
b.p[0] = (__half)a.x;
|
||||
b.p[1] = (__half)a.y;
|
||||
b.x = (__half)a.x;
|
||||
b.y = (__half)a.y;
|
||||
return b;
|
||||
}
|
||||
|
||||
@@ -202,8 +202,8 @@ __device__ __half __float2half(const float a) {
|
||||
|
||||
__device__ __half2 __float2half2_rn(const float a) {
|
||||
__half2 b;
|
||||
b.p[0] = (__half)a;
|
||||
b.p[1] = (__half)a;
|
||||
b.x = (__half)a;
|
||||
b.y = (__half)a;
|
||||
return b;
|
||||
}
|
||||
|
||||
@@ -225,15 +225,15 @@ __device__ __half __float2half_rz(const float a) {
|
||||
|
||||
__device__ __half2 __floats2half2_rn(const float a, const float b) {
|
||||
__half2 c;
|
||||
c.p[0] = (__half)a;
|
||||
c.p[1] = (__half)b;
|
||||
c.x = (__half)a;
|
||||
c.y = (__half)b;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ float2 __half22float2(const __half2 a) {
|
||||
float2 b;
|
||||
b.x = (float)a.p[0];
|
||||
b.y = (float)a.p[1];
|
||||
b.x = (float)a.x;
|
||||
b.y = (float)a.y;
|
||||
return b;
|
||||
}
|
||||
|
||||
@@ -243,8 +243,8 @@ __device__ float __half2float(const __half a) {
|
||||
|
||||
__device__ __half2 half2half2(const __half a) {
|
||||
__half2 b;
|
||||
b.p[0] = a;
|
||||
b.p[1] = a;
|
||||
b.x = a;
|
||||
b.y = a;
|
||||
return b;
|
||||
}
|
||||
|
||||
@@ -358,30 +358,30 @@ __device__ unsigned short int __half_as_ushort(const __half h) {
|
||||
|
||||
__device__ __half2 __halves2half2(const __half a, const __half b) {
|
||||
__half2 c;
|
||||
c.p[0] = a;
|
||||
c.p[1] = b;
|
||||
c.x = a;
|
||||
c.y = b;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ float __high2float(const __half2 a) {
|
||||
return (float)a.p[1];
|
||||
return (float)a.y;
|
||||
}
|
||||
|
||||
__device__ __half __high2half(const __half2 a) {
|
||||
return a.p[1];
|
||||
return a.y;
|
||||
}
|
||||
|
||||
__device__ __half2 __high2half2(const __half2 a) {
|
||||
__half2 b;
|
||||
b.p[0] = a.p[1];
|
||||
b.p[1] = a.p[1];
|
||||
b.x = a.y;
|
||||
b.y = a.y;
|
||||
return b;
|
||||
}
|
||||
|
||||
__device__ __half2 __highs2half2(const __half2 a, const __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = a.p[1];
|
||||
c.p[1] = b.p[1];
|
||||
c.x = a.y;
|
||||
c.y = b.y;
|
||||
return c;
|
||||
}
|
||||
|
||||
@@ -418,38 +418,38 @@ __device__ __half __ll2half_rz(long long int i){
|
||||
}
|
||||
|
||||
__device__ float __low2float(const __half2 a) {
|
||||
return (float)a.p[0];
|
||||
return (float)a.x;
|
||||
}
|
||||
|
||||
__device__ __half __low2half(const __half2 a) {
|
||||
return a.p[0];
|
||||
return a.x;
|
||||
}
|
||||
|
||||
__device__ __half2 __low2half2(const __half2 a, const __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = a.p[0];
|
||||
c.p[1] = b.p[0];
|
||||
c.x = a.x;
|
||||
c.y = b.x;
|
||||
return c;
|
||||
}
|
||||
|
||||
__device__ __half2 __low2half2(const __half2 a) {
|
||||
__half2 b;
|
||||
b.p[0] = a.p[0];
|
||||
b.p[1] = a.p[0];
|
||||
b.x = a.x;
|
||||
b.y = a.x;
|
||||
return b;
|
||||
}
|
||||
|
||||
__device__ __half2 __lowhigh2highlow(const __half2 a) {
|
||||
__half2 b;
|
||||
b.p[0] = a.p[1];
|
||||
b.p[1] = a.p[0];
|
||||
b.x = a.y;
|
||||
b.y = a.x;
|
||||
return b;
|
||||
}
|
||||
|
||||
__device__ __half2 __lows2half2(const __half2 a, const __half2 b) {
|
||||
__half2 c;
|
||||
c.p[0] = a.p[0];
|
||||
c.p[1] = b.p[0];
|
||||
c.y = a.x;
|
||||
c.y = b.x;
|
||||
return c;
|
||||
}
|
||||
|
||||
@@ -542,346 +542,4 @@ typedef struct{
|
||||
};
|
||||
} struct_float;
|
||||
|
||||
#if __clang_major__ == 3
|
||||
|
||||
static __device__ float cvt_half_to_float(__half a){
|
||||
struct_float ret = {0};
|
||||
if(a.x == 0){
|
||||
return 0.0f;
|
||||
}
|
||||
if(a.x == 0x8000){
|
||||
return -0.0f;
|
||||
}
|
||||
ret.u = ((a.x&0x8000)<<16) | (((a.x&0x7c00)+0x1C000)<<13) | ((a.x&0x03FF)<<13);
|
||||
return ret.f;
|
||||
}
|
||||
|
||||
static __device__ __half cvt_float_to_half(float b){
|
||||
struct_float f = {0};
|
||||
__half ret = {0};
|
||||
f.f = b;
|
||||
if(f.f == 0.0f){
|
||||
ret.x = 0;
|
||||
return ret;
|
||||
}
|
||||
if(f.f == -0.0f){
|
||||
ret.x = 0x8000;
|
||||
return ret;
|
||||
}
|
||||
ret.x = ((f.u>>16)&0x8000)|((((f.u&0x7f800000)-0x38000000)>>13)&0x7c00)|((f.u>>13)&0x03ff);
|
||||
return ret;
|
||||
}
|
||||
|
||||
|
||||
__device__ __half __soft_hadd(const __half a, const __half b){
|
||||
return cvt_float_to_half(cvt_half_to_float(a)+cvt_half_to_float(b));
|
||||
}
|
||||
|
||||
__device__ __half __soft_hadd_sat(const __half a, const __half b){
|
||||
float f = cvt_half_to_float(a) + cvt_half_to_float(b);
|
||||
return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f)));
|
||||
}
|
||||
|
||||
__device__ __half __soft_hfma(const __half a, const __half b, const __half c){
|
||||
return cvt_float_to_half(fmaf(cvt_half_to_float(a), cvt_half_to_float(b), cvt_half_to_float(c)));
|
||||
}
|
||||
|
||||
__device__ __half __soft_hfma_sat(const __half a, const __half b, const __half c){
|
||||
float f = fmaf(cvt_half_to_float(a), cvt_half_to_float(b), cvt_half_to_float(c));
|
||||
return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f)));
|
||||
}
|
||||
|
||||
__device__ __half __soft_hmul(const __half a, const __half b){
|
||||
return cvt_float_to_half(cvt_half_to_float(a)*cvt_half_to_float(b));
|
||||
}
|
||||
|
||||
__device__ __half __soft_hmul_sat(const __half a, const __half b){
|
||||
float f = cvt_half_to_float(a) * cvt_half_to_float(b);
|
||||
return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f)));
|
||||
}
|
||||
|
||||
__device__ __half __soft_hneq(const __half a){
|
||||
__half ret = {a.x};
|
||||
ret.x ^= 1 << 15;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half __soft_hsub(const __half a, const __half b){
|
||||
return cvt_float_to_half(cvt_half_to_float(a)-cvt_half_to_float(b));
|
||||
}
|
||||
|
||||
__device__ __half __soft_hsub_sat(const __half a, const __half b){
|
||||
float f = cvt_half_to_float(a) - cvt_half_to_float(b);
|
||||
return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f)));
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
Half2 Arithmetic Instructions
|
||||
*/
|
||||
|
||||
__device__ __half2 __soft_hadd2(const __half2 a, const __half2 b){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hadd(a.p[1], b.p[1]);
|
||||
ret.p[0] = __soft_hadd(a.p[0], b.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hadd2_sat(const __half2 a, const __half2 b){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hadd_sat(a.p[1], b.p[1]);
|
||||
ret.p[0] = __soft_hadd_sat(a.p[0], b.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hfma2(const __half2 a, const __half2 b, const __half2 c){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hfma(a.p[1], b.p[1], c.p[1]);
|
||||
ret.p[0] = __soft_hfma(a.p[0], b.p[0], c.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hfma2_sat(const __half2 a, const __half2 b, const __half2 c){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hfma_sat(a.p[1], b.p[1], c.p[1]);
|
||||
ret.p[0] = __soft_hfma_sat(a.p[0], b.p[0], c.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hmul2(const __half2 a, const __half2 b){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hmul(a.p[1], b.p[1]);
|
||||
ret.p[0] = __soft_hmul(a.p[0], b.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hmul2_sat(const __half2 a, const __half2 b){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hmul_sat(a.p[1], b.p[1]);
|
||||
ret.p[0] = __soft_hmul_sat(a.p[0], b.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hneq2(const __half2 a){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hneq(a.p[1]);
|
||||
ret.p[0] = __soft_hneq(a.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hsub2(const __half2 a, const __half2 b){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hsub(a.p[1], b.p[1]);
|
||||
ret.p[0] = __soft_hsub(a.p[0], b.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hsub2_sat(const __half2 a, const __half2 b){
|
||||
__half2 ret;
|
||||
ret.p[1] = __soft_hsub_sat(a.p[1], b.p[1]);
|
||||
ret.p[0] = __soft_hsub_sat(a.p[0], b.p[0]);
|
||||
return ret;
|
||||
}
|
||||
|
||||
/*
|
||||
Half Cmps
|
||||
*/
|
||||
|
||||
__device__ bool __soft_heq(const __half a, const __half b){
|
||||
return (a.x == b.x ? true:false);
|
||||
}
|
||||
|
||||
__device__ bool __soft_hge(const __half a, const __half b){
|
||||
return (cvt_half_to_float(a) >= cvt_half_to_float(b));
|
||||
}
|
||||
|
||||
__device__ bool __soft_hgt(const __half a, const __half b){
|
||||
return (cvt_half_to_float(a) > cvt_half_to_float(b));
|
||||
}
|
||||
|
||||
__device__ bool __soft_hisinf(const __half a){
|
||||
return ((a.x == __half_neg_inf) ? -1 : (a.x == __half_pos_inf) ? 1 : 0);
|
||||
}
|
||||
|
||||
__device__ bool __soft_hisnan(const __half a){
|
||||
if(((a.x & __half_pos_inf) == a.x) || ((a.x & __half_neg_inf) == a.x)){
|
||||
return true;
|
||||
}else{
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
__device__ bool __soft_hle(const __half a, const __half b){
|
||||
return (cvt_half_to_float(a) <= cvt_half_to_float(b));
|
||||
}
|
||||
|
||||
__device__ bool __soft_hlt(const __half a, const __half b){
|
||||
return (cvt_half_to_float(a) < cvt_half_to_float(b));
|
||||
}
|
||||
|
||||
__device__ bool __soft_hne(const __half a, const __half b){
|
||||
return a.x == b.x ? false : true;
|
||||
}
|
||||
|
||||
/*
|
||||
Half2 Cmps
|
||||
*/
|
||||
|
||||
__device__ bool __soft_hbeq2(const __half2 a, const __half2 b){
|
||||
return __soft_heq(a.p[1], b.p[1]) && __soft_heq(a.p[0], b.p[0]);
|
||||
}
|
||||
|
||||
__device__ bool __soft_hbge2(const __half2 a, const __half2 b){
|
||||
return __soft_hge(a.p[1], b.p[1]) && __soft_hge(a.p[0], b.p[0]);
|
||||
}
|
||||
|
||||
__device__ bool __soft_hbgt2(const __half2 a, const __half2 b){
|
||||
return __soft_hgt(a.p[1], b.p[1]) && __soft_hgt(a.p[0], b.p[0]);
|
||||
}
|
||||
|
||||
__device__ bool __soft_hble2(const __half2 a, const __half2 b){
|
||||
return __soft_hle(a.p[1], b.p[1]) && __soft_hle(a.p[0], b.p[0]);
|
||||
}
|
||||
|
||||
__device__ bool __soft_hblt2(const __half2 a, const __half2 b){
|
||||
return __soft_hlt(a.p[1], b.p[1]) && __soft_hlt(a.p[0], b.p[0]);
|
||||
}
|
||||
|
||||
__device__ bool __soft_hbne2(const __half2 a, const __half2 b){
|
||||
return __soft_hne(a.p[1], b.p[1]) && __soft_hne(a.p[0], b.p[0]);
|
||||
}
|
||||
|
||||
|
||||
|
||||
__device__ __half2 __soft_heq2(const __half2 a, const __half2 b){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = (__soft_heq(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float;
|
||||
ret.p[0] = (__soft_heq(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hge2(const __half2 a, const __half2 b){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = (__soft_hge(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float;
|
||||
ret.p[0] = (__soft_hge(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hgt2(const __half2 a, const __half2 b){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = (__soft_hgt(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float;
|
||||
ret.p[0] = (__soft_hgt(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hisnan2(const __half2 a){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = __soft_hisnan(a.p[1]) ? __half_value_one_float : __half_value_zero_float;
|
||||
ret.p[0] = __soft_hisnan(a.p[0]) ? __half_value_one_float : __half_value_zero_float;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hle2(const __half2 a, const __half2 b){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = (__soft_hle(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float;
|
||||
ret.p[0] = (__soft_hle(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hlt2(const __half2 a, const __half2 b){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = (__soft_hlt(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float;
|
||||
ret.p[0] = (__soft_hlt(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_hne2(const __half2 a, const __half2 b){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = (__soft_hne(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float;
|
||||
ret.p[0] = (__soft_hne(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float;
|
||||
return ret;
|
||||
}
|
||||
|
||||
/*
|
||||
Half Cnvs and Data Mvmnt
|
||||
*/
|
||||
|
||||
__device__ __half2 __soft_float22half2_rn(const float2 a){
|
||||
__half2 ret = {0};
|
||||
ret.p[1] = cvt_float_to_half(a.x);
|
||||
ret.p[0] = cvt_float_to_half(a.y);
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ __half __soft_float2half(const float a){
|
||||
return cvt_float_to_half(a);
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_float2half2_rn(const float a){
|
||||
__half ret = cvt_float_to_half(a);
|
||||
return {ret, ret};
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_floats2half2_rn(const float a, const float b){
|
||||
return {cvt_float_to_half(a), cvt_float_to_half(b)};
|
||||
}
|
||||
|
||||
__device__ float2 __soft_half22float2(const __half2 a){
|
||||
return {cvt_half_to_float(a.p[1]), cvt_half_to_float(a.p[0])};
|
||||
}
|
||||
|
||||
__device__ float __soft_half2float(const __half a){
|
||||
return cvt_half_to_float(a);
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_half2half2(const __half a){
|
||||
return {a,a};
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_halves2half2(const __half a, const __half b){
|
||||
return {a,b};
|
||||
}
|
||||
|
||||
__device__ float __soft_high2float(const __half2 a){
|
||||
return cvt_half_to_float(a.p[1]);
|
||||
}
|
||||
|
||||
__device__ __half __soft_high2half(const __half2 a){
|
||||
return a.p[1];
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_high2half2(const __half2 a){
|
||||
return {a.p[1], a.p[1]};
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_highs2half2(const __half2 a, const __half2 b){
|
||||
return {a.p[1], b.p[1]};
|
||||
}
|
||||
|
||||
__device__ float __soft_low2float(const __half2 a){
|
||||
return cvt_half_to_float(a.p[0]);
|
||||
}
|
||||
|
||||
__device__ __half __soft_low2half(const __half2 a){
|
||||
return a.p[0];
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_low2half2(const __half2 a){
|
||||
return {a.p[0], a.p[0]};
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_lows2half2(const __half2 a, const __half2 b){
|
||||
return {a.p[0], b.p[0]};
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_lowhigh2highlow(const __half2 a){
|
||||
return {a.p[0], a.p[1]};
|
||||
}
|
||||
|
||||
__device__ __half2 __soft_low2half2(const __half2 a, const __half2 b){
|
||||
return {a.p[0], b.p[0]};
|
||||
}
|
||||
|
||||
|
||||
|
||||
#endif
|
||||
|
||||
@@ -2,89 +2,122 @@ target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:
|
||||
target triple = "amdgcn--amdhsa"
|
||||
|
||||
|
||||
define i32 @__hip_hc_ir_hadd2_int(i32 %a, i32 %b) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_add_f16 $0, $1, $2","=v,v,v"(i32 %a, i32 %b)
|
||||
tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_hadd2_int(<2 x half> %a, <2 x half> %b) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = bitcast <2 x half> %b to i32
|
||||
%3 = tail call i32 asm sideeffect "v_add_f16 $0, $1, $2","=v,v,v"(i32 %1, i32 %2)
|
||||
tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %3, i32 %1, i32 %2)
|
||||
%4 = bitcast i32 %3 to <2 x half>
|
||||
ret <2 x half> %4
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_hfma2_int(i32 %a, i32 %b, i32 %c) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_mad_f16 $0, $1, $2, $3","=v,v,v,v"(i32 %a, i32 %b, i32 %c)
|
||||
tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b)
|
||||
tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %1, i32 %c)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_hfma2_int(<2 x half> %a, <2 x half> %b, <2 x half> %c) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = bitcast <2 x half> %b to i32
|
||||
%3 = bitcast <2 x half> %c to i32
|
||||
%4 = tail call i32 asm sideeffect "v_mad_f16 $0, $1, $2, $3","=v,v,v,v"(i32 %1, i32 %2, i32 %3)
|
||||
tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %4, i32 %1, i32 %2)
|
||||
tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %4, i32 %4, i32 %3)
|
||||
%5 = bitcast i32 %4 to <2 x half>
|
||||
ret <2 x half> %5
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_hmul2_int(i32 %a, i32 %b) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_mul_f16 $0, $1, $2","=v,v,v"(i32 %a, i32 %b)
|
||||
tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_hmul2_int(<2 x half> %a, <2 x half> %b) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = bitcast <2 x half> %b to i32
|
||||
%3 = tail call i32 asm sideeffect "v_mul_f16 $0, $1, $2","=v,v,v"(i32 %1, i32 %2)
|
||||
tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %3, i32 %1, i32 %2)
|
||||
%4 = bitcast i32 %3 to <2 x half>
|
||||
ret <2 x half> %4
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_hsub2_int(i32 %a, i32 %b) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_sub_f16 $0, $1, $2","=v,v,v"(i32 %a, i32 %b)
|
||||
tail call void asm sideeffect "v_sub_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_hsub2_int(<2 x half> %a, <2 x half> %b) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = bitcast <2 x half> %b to i32
|
||||
%3 = tail call i32 asm sideeffect "v_sub_f16 $0, $1, $2","=v,v,v"(i32 %1, i32 %2)
|
||||
tail call void asm sideeffect "v_sub_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %3, i32 %1, i32 %2)
|
||||
%4 = bitcast i32 %3 to <2 x half>
|
||||
ret <2 x half> %4
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2ceil_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_ceil_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_ceil_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2ceil_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_ceil_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_ceil_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2cos_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_cos_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_cos_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2cos_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_cos_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_cos_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2exp2_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_exp_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_exp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2exp2_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_exp_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_exp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2floor_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_floor_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_floor_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2floor_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_floor_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_floor_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2log2_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_log_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_log_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2log2_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_log_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_log_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2rcp_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_rcp_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_rcp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2rcp_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_rcp_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_rcp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2rsqrt_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_rsq_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_rsq_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2rsqrt_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_rsq_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_rsq_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2sin_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_sin_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_sin_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2sin_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_sin_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_sin_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2sqrt_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_sqrt_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_sqrt_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2sqrt_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_sqrt_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_sqrt_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
define i32 @__hip_hc_ir_h2trunc_int(i32 %a) #1 {
|
||||
%1 = tail call i32 asm sideeffect "v_trunc_f16 $0, $1","=v,v"(i32 %a)
|
||||
tail call void asm sideeffect "v_trunc_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
|
||||
ret i32 %1
|
||||
define <2 x half> @__hip_hc_ir_h2trunc_int(<2 x half> %a) #1 {
|
||||
%1 = bitcast <2 x half> %a to i32
|
||||
%2 = tail call i32 asm sideeffect "v_trunc_f16 $0, $1","=v,v"(i32 %1)
|
||||
tail call void asm sideeffect "v_trunc_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1)
|
||||
%3 = bitcast i32 %2 to <2 x half>
|
||||
ret <2 x half> %3
|
||||
}
|
||||
|
||||
attributes #1 = { alwaysinline nounwind }
|
||||
|
||||
Ссылка в новой задаче
Block a user