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

Guard the OCML rounded operations instead

Browse Source 
Instead of commenting all these functions out, guard the functions with a macro OCML_BASIC_ROUNDED_OPERATIONS.
This commit is contained in:
Aaron Enye Shi
2018-11-06 16:32:14 +00:00
parent 7223277aa3
commit e59c33250a
3 changed files with 236 additions and 216 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
+28 -24
View File
@@ -34,34 +34,38 @@ THE SOFTWARE.
#pragma clang diagnostic ignored "-Wunused-variable"
__device__ void double_precision_intrinsics() {
// __dadd_rd(0.0, 1.0);
// __dadd_rn(0.0, 1.0);
// __dadd_ru(0.0, 1.0);
// __dadd_rz(0.0, 1.0);
// __ddiv_rd(0.0, 1.0);
// __ddiv_rn(0.0, 1.0);
// __ddiv_ru(0.0, 1.0);
// __ddiv_rz(0.0, 1.0);
// __dmul_rd(1.0, 2.0);
// __dmul_rn(1.0, 2.0);
// __dmul_ru(1.0, 2.0);
// __dmul_rz(1.0, 2.0);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__dadd_rd(0.0, 1.0);
__dadd_rn(0.0, 1.0);
__dadd_ru(0.0, 1.0);
__dadd_rz(0.0, 1.0);
__ddiv_rd(0.0, 1.0);
__ddiv_rn(0.0, 1.0);
__ddiv_ru(0.0, 1.0);
__ddiv_rz(0.0, 1.0);
__dmul_rd(1.0, 2.0);
__dmul_rn(1.0, 2.0);
__dmul_ru(1.0, 2.0);
__dmul_rz(1.0, 2.0);
#endif
__drcp_rd(2.0);
__drcp_rn(2.0);
__drcp_ru(2.0);
__drcp_rz(2.0);
// __dsqrt_rd(4.0);
// __dsqrt_rn(4.0);
// __dsqrt_ru(4.0);
// __dsqrt_rz(4.0);
// __dsub_rd(2.0, 1.0);
// __dsub_rn(2.0, 1.0);
// __dsub_ru(2.0, 1.0);
// __dsub_rz(2.0, 1.0);
// __fma_rd(1.0, 2.0, 3.0);
// __fma_rn(1.0, 2.0, 3.0);
// __fma_ru(1.0, 2.0, 3.0);
// __fma_rz(1.0, 2.0, 3.0);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__dsqrt_rd(4.0);
__dsqrt_rn(4.0);
__dsqrt_ru(4.0);
__dsqrt_rz(4.0);
__dsub_rd(2.0, 1.0);
__dsub_rn(2.0, 1.0);
__dsub_ru(2.0, 1.0);
__dsub_rz(2.0, 1.0);
__fma_rd(1.0, 2.0, 3.0);
__fma_rn(1.0, 2.0, 3.0);
__fma_ru(1.0, 2.0, 3.0);
__fma_rz(1.0, 2.0, 3.0);
#endif
}
__global__ void compileDoublePrecisionIntrinsics(int ignored) {
tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
+30 -24
View File
@@ -39,36 +39,42 @@ __device__ void single_precision_intrinsics() {
__cosf(0.0f);
__exp10f(0.0f);
__expf(0.0f);
// __fadd_rd(0.0f, 1.0f);
// __fadd_rn(0.0f, 1.0f);
// __fadd_ru(0.0f, 1.0f);
// __fadd_rz(0.0f, 1.0f);
// __fdiv_rd(4.0f, 2.0f);
// __fdiv_rn(4.0f, 2.0f);
// __fdiv_ru(4.0f, 2.0f);
// __fdiv_rz(4.0f, 2.0f);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__fadd_rd(0.0f, 1.0f);
__fadd_rn(0.0f, 1.0f);
__fadd_ru(0.0f, 1.0f);
__fadd_rz(0.0f, 1.0f);
__fdiv_rd(4.0f, 2.0f);
__fdiv_rn(4.0f, 2.0f);
__fdiv_ru(4.0f, 2.0f);
__fdiv_rz(4.0f, 2.0f);
#endif
__fdividef(4.0f, 2.0f);
// __fmaf_rd(1.0f, 2.0f, 3.0f);
// __fmaf_rn(1.0f, 2.0f, 3.0f);
// __fmaf_ru(1.0f, 2.0f, 3.0f);
// __fmaf_rz(1.0f, 2.0f, 3.0f);
// __fmul_rd(1.0f, 2.0f);
// __fmul_rn(1.0f, 2.0f);
// __fmul_ru(1.0f, 2.0f);
// __fmul_rz(1.0f, 2.0f);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__fmaf_rd(1.0f, 2.0f, 3.0f);
__fmaf_rn(1.0f, 2.0f, 3.0f);
__fmaf_ru(1.0f, 2.0f, 3.0f);
__fmaf_rz(1.0f, 2.0f, 3.0f);
__fmul_rd(1.0f, 2.0f);
__fmul_rn(1.0f, 2.0f);
__fmul_ru(1.0f, 2.0f);
__fmul_rz(1.0f, 2.0f);
#endif
__frcp_rd(2.0f);
__frcp_rn(2.0f);
__frcp_ru(2.0f);
__frcp_rz(2.0f);
__frsqrt_rn(4.0f);
// __fsqrt_rd(4.0f);
// __fsqrt_rn(4.0f);
// __fsqrt_ru(4.0f);
// __fsqrt_rz(4.0f);
// __fsub_rd(2.0f, 1.0f);
// __fsub_rn(2.0f, 1.0f);
// __fsub_ru(2.0f, 1.0f);
// __fsub_rz(2.0f, 1.0f);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__fsqrt_rd(4.0f);
__fsqrt_rn(4.0f);
__fsqrt_ru(4.0f);
__fsqrt_rz(4.0f);
__fsub_rd(2.0f, 1.0f);
__fsub_rn(2.0f, 1.0f);
__fsub_ru(2.0f, 1.0f);
__fsub_rz(2.0f, 1.0f);
#endif
__log10f(1.0f);
__log2f(1.0f);
__logf(1.0f);