AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

added half2 math operations

Explorar el Código 
1. They use SDWA + LLVM IR
2. Added these functions to test
3. Need to do exp, exp10, log, log10, rint

Change-Id: I06176acc6cb8bb054495310531777406a41b54e4
Este commit está contenido en:
Aditya Atluri
2017-01-13 12:27:11 -06:00
padre fe38e9652b
commit eff68c989a
Se han modificado 3 ficheros con 152 adiciones y 0 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
hipamd/include/hip/hcc_detail/hip_fp16.h
+76
Ver fichero
@@ -69,6 +69,17 @@ 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);
__device__ static inline __half __hadd(const __half a, const __half b) {
return __hip_hc_ir_hadd_half(a, b);
}
@@ -682,6 +693,71 @@ __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 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 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
hipamd/src/hip_ir.ll
+60
Ver fichero
@@ -86,4 +86,64 @@ define i32 @__hip_hc_ir_hsub2_int(i32 %a, i32 %b) #1 {
ret i32 %1
}
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 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 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 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 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 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 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 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 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 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
}
attributes #1 = { alwaysinline nounwind }
hipamd/tests/src/deviceLib/hipTestHalf.cpp
+16
Ver fichero
@@ -88,6 +88,21 @@ __global__ void CheckHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2
Out[7] = __hsub2(In1[7], In2[7]);
Out[8] = __hsub2_sat(In1[8], In2[8]);
Out[9] = h2div(In1[9], In2[9]);
Out[10] = h2ceil(In1[10]);
Out[11] = h2cos(In1[11]);
// Out[12] = h2exp(In1[12]);
// Out[13] = h2exp10(In1[13]);
Out[14] = h2exp2(In1[14]);
Out[15] = h2floor(In1[15]);
// Out[16] = h2log(In1[16]);
// Out[17] = h2log10(In1[17]);
Out[18] = h2log2(In1[18]);
Out[19] = h2rcp(In1[19]);
// Out[20] = h2rint(In1[20]);
Out[21] = h2rsqrt(In1[21]);
Out[22] = h2sin(In1[22]);
Out[23] = h2sqrt(In1[23]);
Out[24] = h2trunc(In1[24]);
}
__global__ void CheckCmpHalf(hipLaunchParm lp, __half* In1, __half* In2, bool* Out) {
@@ -109,6 +124,7 @@ __global__ void CheckCmpHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __ha
Out[5] = __hle2(In1[5], In2[5]);
Out[6] = __hlt2(In1[6], In2[6]);
Out[7] = __hne2(In1[7], In2[7]);
}
int main(){