From e9daf7624b2d61fef8634c8b43b016950d8df29f Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Tue, 23 Oct 2018 21:13:11 +0800 Subject: [PATCH 1/9] Make correct checking of the returned hipDeviceptr_t from read_global_description() [ROCm/hip commit: de5f47a98410e08761199be33134fc012e208353] --- projects/hip/src/hip_module.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index a6d486b6de..bb419740e3 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -342,7 +342,7 @@ hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, hi tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), it0->second.cend(), name); - return dptr ? hipSuccess : hipErrorNotFound; + return *dptr ? hipSuccess : hipErrorNotFound; } hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, const char* name) { @@ -367,7 +367,7 @@ hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, c tie(*dptr, *bytes) = read_global_description(it->second.cbegin(), it->second.cend(), name); - return dptr ? hipSuccess : hipErrorNotFound; + return *dptr ? hipSuccess : hipErrorNotFound; } hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname) { From 3e91f67882636b56106a795b17a99ab136639eb6 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 25 Oct 2018 12:19:32 -0400 Subject: [PATCH 2/9] Adding checks and debug output for fat binary for hip-clang [ROCm/hip commit: 062398c72f81940c71f0f7728ebf58feaf790ef3] --- projects/hip/src/hip_clang.cpp | 53 ++++++++++++++++++++++++++-------- 1 file changed, 41 insertions(+), 12 deletions(-) diff --git a/projects/hip/src/hip_clang.cpp b/projects/hip/src/hip_clang.cpp index 15a96d298a..6067edd289 100644 --- a/projects/hip/src/hip_clang.cpp +++ b/projects/hip/src/hip_clang.cpp @@ -86,6 +86,7 @@ __hipRegisterFatBinary(const void* data) std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)], desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)}; + tprintf(DB_FB, "Found bundle for %s\n", target.c_str()); for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) { hsa_agent_t agent = g_allAgents[deviceId + 1]; @@ -110,10 +111,25 @@ __hipRegisterFatBinary(const void* data) if (module->executable.handle) { modules->at(deviceId) = module; + tprintf(DB_FB, "Loaded code object for %s\n", name); + } else { + fprintf(stderr, "Failed to load code object for %s\n", name); + abort(); } } } + for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) { + hsa_agent_t agent = g_allAgents[deviceId + 1]; + + char name[64] = {}; + hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name); + if (!(*modules)[deviceId]) { + fprintf(stderr, "No device code bundle for %s\n", name); + abort(); + } + } + tprintf(DB_FB, "__hipRegisterFatBinary succeeds and returns %p\n", modules); return modules; } @@ -132,13 +148,18 @@ extern "C" void __hipRegisterFunction( dim3* gridDim, int* wSize) { + HIP_INIT_API(modules, hostFunction, deviceFunction, deviceName); std::vector functions{g_deviceCnt}; + assert(modules && modules->size() >= g_deviceCnt); for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) { hipFunction_t function; if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName)) { functions[deviceId] = function; } + else { + tprintf(DB_FB, "missing kernel %s for device %d\n", deviceName, deviceId); + } } g_functions.insert(std::make_pair(hostFunction, std::move(functions))); @@ -180,6 +201,7 @@ hipError_t hipSetupArgument( size_t size, size_t offset) { + HIP_INIT_API(arg, size, offset); auto ctx = ihipGetTlsDefaultCtx(); LockedAccessor_CtxCrit_t crit(ctx->criticalData()); auto& arguments = crit->_execStack.top()._arguments; @@ -194,6 +216,7 @@ hipError_t hipSetupArgument( hipError_t hipLaunchByPtr(const void *hostFunction) { + HIP_INIT_API(hostFunction); ihipExec_t exec; { auto ctx = ihipGetTlsDefaultCtx(); @@ -213,20 +236,26 @@ hipError_t hipLaunchByPtr(const void *hostFunction) deviceId = 0; } + hipError_t e = hipSuccess; decltype(g_functions)::iterator it; - if ((it = g_functions.find(hostFunction)) == g_functions.end()) - return hipErrorUnknown; + if ((it = g_functions.find(hostFunction)) == g_functions.end()) { + e = hipErrorUnknown; + fprintf(stderr, "kernel %p not found!\n", hostFunction); + abort(); + } else { + size_t size = exec._arguments.size(); + void *extra[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0], + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END + }; - size_t size = exec._arguments.size(); - void *extra[] = { - HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0], - HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, - HIP_LAUNCH_PARAM_END - }; + e = hipModuleLaunchKernel(it->second[deviceId], + exec._gridDim.x, exec._gridDim.y, exec._gridDim.z, + exec._blockDim.x, exec._blockDim.y, exec._blockDim.z, + exec._sharedMem, exec._hStream, nullptr, extra); + } - return hipModuleLaunchKernel(it->second[deviceId], - exec._gridDim.x, exec._gridDim.y, exec._gridDim.z, - exec._blockDim.x, exec._blockDim.y, exec._blockDim.z, - exec._sharedMem, exec._hStream, nullptr, extra); + return ihipLogStatus(e); } From fa429022e10f02c6a3ff36f6ef76c5be34d7dc3e Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 26 Oct 2018 14:11:18 -0400 Subject: [PATCH 3/9] Add HIP_DUMP_CODE_OBJECT [ROCm/hip commit: f5d8842f6a74d7c01de88064b56b230ba46a597a] --- projects/hip/src/hip_clang.cpp | 11 +++++++++++ projects/hip/src/hip_hcc.cpp | 6 ++++++ projects/hip/src/hip_hcc_internal.h | 2 +- 3 files changed, 18 insertions(+), 1 deletion(-) diff --git a/projects/hip/src/hip_clang.cpp b/projects/hip/src/hip_clang.cpp index 6067edd289..cfd75df562 100644 --- a/projects/hip/src/hip_clang.cpp +++ b/projects/hip/src/hip_clang.cpp @@ -22,6 +22,7 @@ THE SOFTWARE. #include #include +#include #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" @@ -112,6 +113,16 @@ __hipRegisterFatBinary(const void* data) if (module->executable.handle) { modules->at(deviceId) = module; tprintf(DB_FB, "Loaded code object for %s\n", name); + if (HIP_DUMP_CODE_OBJECT) { + char fname[30]; + static std::atomic index; + sprintf(fname, "__hip_dump_code_object%04d.o", index++); + tprintf(DB_FB, "Dump code object %s\n", fname); + std::ofstream ofs; + ofs.open(fname, std::ios::binary); + ofs << image; + ofs.close(); + } } else { fprintf(stderr, "Failed to load code object for %s\n", name); abort(); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index eff93da847..e152e7ba69 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -97,6 +97,8 @@ int HIP_INIT_ALLOC = -1; int HIP_SYNC_STREAM_WAIT = 0; int HIP_FORCE_NULL_STREAM = 0; +int HIP_DUMP_CODE_OBJECT = 0; + #if (__hcc_workweek__ >= 17300) // Make sure we have required bug fix in HCC @@ -1294,6 +1296,10 @@ void HipReadEnv() { "overridden by specifying hipEventReleaseToSystem or hipEventReleaseToDevice flag " "when creating the event."); + READ_ENV_I(release, HIP_DUMP_CODE_OBJECT, 0, + "If set, dump code object as __hip_dump_code_object[nnnn].o in the current directory," + "where nnnn is the index number."); + // Some flags have both compile-time and runtime flags - generate a warning if user enables the // runtime flag but the compile-time flag is disabled. if (HIP_DB && !COMPILE_HIP_DB) { diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index d64a4a4cbe..8102f066de 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -83,11 +83,11 @@ extern int HIP_SYNC_NULL_STREAM; extern int HIP_INIT_ALLOC; extern int HIP_FORCE_NULL_STREAM; +extern int HIP_DUMP_CODE_OBJECT; // TODO - remove when this is standard behavior. extern int HCC_OPT_FLUSH; - // Class to assign a short TID to each new thread, for HIP debugging purposes. class TidInfo { public: From d4815f4178f1755186d69dfb562cd31ce8892dbf Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 31 Oct 2018 03:22:38 +0530 Subject: [PATCH 4/9] Fixes global symbols tracking in hip_module [ROCm/hip commit: b270313129be98995fca63e32e37ea796fd62a98] --- projects/hip/src/hip_module.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index a6d486b6de..0aa171631c 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -267,7 +267,12 @@ inline void track(const Agent_global& x) { hc::AmPointerInfo ptr_info(nullptr, x.address, x.address, x.byte_cnt, device->_acc, true, false); hc::am_memtracker_add(x.address, ptr_info); +#if USE_APP_PTR_FOR_CTX + hc::am_memtracker_update(x.address, device->_deviceId, 0u, ihipGetTlsDefaultCtx()); +#else hc::am_memtracker_update(x.address, device->_deviceId, 0u); +#endif + } template > From 1159b4aa058ad83ed6a4cf47c9af3d726b341fd6 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 1 Nov 2018 16:20:35 -0400 Subject: [PATCH 5/9] Move the global arrays for hip malloc/free from a header into a source file such that there's only an unique copy in an executable and prevent wasting static memory on the host Change-Id: Id5b62766f77809c8d7b47892cb7149c490dcbdb9 [ROCm/hip commit: 0ff408a56cef3b1dfcd6b622d180bf4ef22de040] --- projects/hip/include/hip/hcc_detail/hip_memory.h | 9 ++++----- projects/hip/src/hip_memory.cpp | 3 +++ 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_memory.h b/projects/hip/include/hip/hcc_detail/hip_memory.h index 739e488ca3..2c9ec1b7c3 100644 --- a/projects/hip/include/hip/hcc_detail/hip_memory.h +++ b/projects/hip/include/hip/hcc_detail/hip_memory.h @@ -27,6 +27,8 @@ THE SOFTWARE. // HIP heap is implemented as a global array with fixed size. Users may define // __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap. +#if __HCC__ || __HIP__ + // Size of page in bytes. #ifndef __HIP_SIZE_OF_PAGE #define __HIP_SIZE_OF_PAGE 64 @@ -39,11 +41,8 @@ THE SOFTWARE. #define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE) -#if __HCC__ || __HIP__ - -__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; -__attribute__((weak)) __device__ - uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; +extern __device__ char __hip_device_heap[]; +extern __device__ uint32_t __hip_device_page_flag[]; extern "C" inline __device__ void* __hip_malloc(size_t size) { char* heap = (char*)__hip_device_heap; diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 93ac527826..4ea5b24f43 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -28,6 +28,9 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" +__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; +__device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; + // Internal HIP APIS: namespace hip_internal { From 1c6f47ef5504aea287e41f4bc8421fa1d73b5e57 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 5 Nov 2018 22:34:16 +0000 Subject: [PATCH 6/9] Remove non-working non-default-rounded math apis In ROCm-Device-Libs, they have dropped the non-default-rounded versions of add, sub, mul, div, sqrt and fma. Therefore, ocml has removed the rte, rtp, rtn, and rtz counterparts. This will remove the same math APIs in HIP for _ru, _rd, _rn, and _rz. [ROCm/hip commit: 2b108a2cfdbc8c1f2e0b8f59271fc7f1224dce80] --- .../include/hip/hcc_detail/math_functions.h | 336 +++++++++--------- .../hip/include/hip/hcc_detail/math_fwd.h | 50 ++- .../hipDoublePrecisionIntrinsics.cpp | 48 +-- .../hipSinglePrecisionIntrinsics.cpp | 48 +-- 4 files changed, 265 insertions(+), 217 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index b12e7aca89..8a6091858b 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -514,69 +514,69 @@ float __exp10f(float x) { return __ocml_exp10_f32(x); } __DEVICE__ inline float __expf(float x) { return __ocml_exp_f32(x); } -__DEVICE__ -inline -float __fadd_rd(float x, float y) { return __ocml_add_rtp_f32(x, y); } -__DEVICE__ -inline -float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); } -__DEVICE__ -inline -float __fadd_ru(float x, float y) { return __ocml_add_rtn_f32(x, y); } -__DEVICE__ -inline -float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } -__DEVICE__ -inline -float __fdiv_rd(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fdiv_rn(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fdiv_ru(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fdiv_rz(float x, float y) { return x / y; } +// __DEVICE__ +// inline +// float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); } +// __DEVICE__ +// inline +// float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); } +// __DEVICE__ +// inline +// float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); } __DEVICE__ inline float __fdividef(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fmaf_rd(float x, float y, float z) -{ - return __ocml_fma_rtp_f32(x, y, z); -} -__DEVICE__ -inline -float __fmaf_rn(float x, float y, float z) -{ - return __ocml_fma_rte_f32(x, y, z); -} -__DEVICE__ -inline -float __fmaf_ru(float x, float y, float z) -{ - return __ocml_fma_rtn_f32(x, y, z); -} -__DEVICE__ -inline -float __fmaf_rz(float x, float y, float z) -{ - return __ocml_fma_rtz_f32(x, y, z); -} -__DEVICE__ -inline -float __fmul_rd(float x, float y) { return __ocml_mul_rtp_f32(x, y); } -__DEVICE__ -inline -float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); } -__DEVICE__ -inline -float __fmul_ru(float x, float y) { return __ocml_mul_rtn_f32(x, y); } -__DEVICE__ -inline -float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); } +// __DEVICE__ +// inline +// float __fmaf_rd(float x, float y, float z) +// { +// return __ocml_fma_rtn_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmaf_rn(float x, float y, float z) +// { +// return __ocml_fma_rte_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmaf_ru(float x, float y, float z) +// { +// return __ocml_fma_rtp_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmaf_rz(float x, float y, float z) +// { +// return __ocml_fma_rtz_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); } +// __DEVICE__ +// inline +// float __fmul_ru(float x, float y) { return __ocml_mul_rtp_f32(x, y); } +// __DEVICE__ +// inline +// float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); } __DEVICE__ inline float __frcp_rd(float x) { return __llvm_amdgcn_rcp_f32(x); } @@ -592,30 +592,30 @@ float __frcp_rz(float x) { return __llvm_amdgcn_rcp_f32(x); } __DEVICE__ inline float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); } -__DEVICE__ -inline -float __fsqrt_rd(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsqrt_ru(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsqrt_rz(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsub_rd(float x, float y) { return __ocml_sub_rtp_f32(x, y); } -__DEVICE__ -inline -float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); } -__DEVICE__ -inline -float __fsub_ru(float x, float y) { return __ocml_sub_rtn_f32(x, y); } -__DEVICE__ -inline -float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); } +// __DEVICE__ +// inline +// float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); } +// __DEVICE__ +// inline +// float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); } +// __DEVICE__ +// inline +// float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); } +// __DEVICE__ +// inline +// float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); } +// __DEVICE__ +// inline +// float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); } +// __DEVICE__ +// inline +// float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); } +// __DEVICE__ +// inline +// float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); } __DEVICE__ inline float __log10f(float x) { return __ocml_log10_f32(x); } @@ -1034,42 +1034,42 @@ double yn(int n, double x) } // BEGIN INTRINSICS -__DEVICE__ -inline -double __dadd_rd(double x, double y) { return __ocml_add_rtp_f64(x, y); } -__DEVICE__ -inline -double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); } -__DEVICE__ -inline -double __dadd_ru(double x, double y) { return __ocml_add_rtn_f64(x, y); } -__DEVICE__ -inline -double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } -__DEVICE__ -inline -double __ddiv_rd(double x, double y) { return x / y; } -__DEVICE__ -inline -double __ddiv_rn(double x, double y) { return x / y; } -__DEVICE__ -inline -double __ddiv_ru(double x, double y) { return x / y; } -__DEVICE__ -inline -double __ddiv_rz(double x, double y) { return x / y; } -__DEVICE__ -inline -double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); } -__DEVICE__ -inline -double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); } -__DEVICE__ -inline -double __dmul_ru(double x, double y) { return __ocml_mul_rtn_f64(x, y); } -__DEVICE__ -inline -double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); } +// __DEVICE__ +// inline +// double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); } +// __DEVICE__ +// inline +// double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); } __DEVICE__ inline double __drcp_rd(double x) { return __llvm_amdgcn_rcp_f64(x); } @@ -1082,54 +1082,54 @@ double __drcp_ru(double x) { return __llvm_amdgcn_rcp_f64(x); } __DEVICE__ inline double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); } -__DEVICE__ -inline -double __dsqrt_rd(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsqrt_rn(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsqrt_ru(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsqrt_rz(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsub_rd(double x, double y) { return __ocml_sub_rtp_f64(x, y); } -__DEVICE__ -inline -double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); } -__DEVICE__ -inline -double __dsub_ru(double x, double y) { return __ocml_sub_rtn_f64(x, y); } -__DEVICE__ -inline -double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); } -__DEVICE__ -inline -double __fma_rd(double x, double y, double z) -{ - return __ocml_fma_rtp_f64(x, y, z); -} -__DEVICE__ -inline -double __fma_rn(double x, double y, double z) -{ - return __ocml_fma_rte_f64(x, y, z); -} -__DEVICE__ -inline -double __fma_ru(double x, double y, double z) -{ - return __ocml_fma_rtn_f64(x, y, z); -} -__DEVICE__ -inline -double __fma_rz(double x, double y, double z) -{ - return __ocml_fma_rtz_f64(x, y, z); -} +// __DEVICE__ +// inline +// double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); } +// __DEVICE__ +// inline +// double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); } +// __DEVICE__ +// inline +// double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); } +// __DEVICE__ +// inline +// double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); } +// __DEVICE__ +// inline +// double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); } +// __DEVICE__ +// inline +// double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); } +// __DEVICE__ +// inline +// double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __fma_rd(double x, double y, double z) +// { +// return __ocml_fma_rtn_f64(x, y, z); +// } +// __DEVICE__ +// inline +// double __fma_rn(double x, double y, double z) +// { +// return __ocml_fma_rte_f64(x, y, z); +// } +// __DEVICE__ +// inline +// double __fma_ru(double x, double y, double z) +// { +// return __ocml_fma_rtp_f64(x, y, z); +// } +// __DEVICE__ +// inline +// double __fma_rz(double x, double y, double z) +// { +// return __ocml_fma_rtz_f64(x, y, z); +// } // END INTRINSICS // END DOUBLE diff --git a/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h index 404c2f81d5..e5594924ba 100644 --- a/projects/hip/include/hip/hcc_detail/math_fwd.h +++ b/projects/hip/include/hip/hcc_detail/math_fwd.h @@ -288,6 +288,30 @@ __attribute__((const)) float __ocml_mul_rtz_f32(float, float); __device__ __attribute__((const)) +float __ocml_div_rte_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_div_rtn_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_div_rtp_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_div_rtz_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rte_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rtn_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rtp_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rtz_f32(float, float); +__device__ +__attribute__((const)) float __ocml_fma_rte_f32(float, float, float); __device__ __attribute__((const)) @@ -572,6 +596,30 @@ __attribute__((const)) double __ocml_mul_rtz_f64(double, double); __device__ __attribute__((const)) +double __ocml_div_rte_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_div_rtn_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_div_rtp_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_div_rtz_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rte_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rtn_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rtp_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rtz_f64(double, double); +__device__ +__attribute__((const)) double __ocml_fma_rte_f64(double, double, double); __device__ __attribute__((const)) @@ -594,4 +642,4 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); #if defined(__cplusplus) } // extern "C" -#endif \ No newline at end of file +#endif diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 17cd82c9ab..f6c515c03a 100644 --- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -34,34 +34,34 @@ 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); + // __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); __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); + // __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); } __global__ void compileDoublePrecisionIntrinsics(int ignored) { diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index f3d2a36931..623ea08a94 100644 --- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp @@ -39,36 +39,36 @@ __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); + // __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); __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); + // __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); __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); + // __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); __log10f(1.0f); __log2f(1.0f); __logf(1.0f); From 914506055ee9223d4c670e6a91ee168dc2c5e8c8 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 5 Nov 2018 22:54:55 +0000 Subject: [PATCH 7/9] Update hip-math-api doc to remove non-default-rounded [ROCm/hip commit: 7223277aa35544b89fceeb66cd94b16e8ea82156] --- projects/hip/docs/markdown/hip-math-api.md | 102 +++++++++++---------- 1 file changed, 55 insertions(+), 47 deletions(-) diff --git a/projects/hip/docs/markdown/hip-math-api.md b/projects/hip/docs/markdown/hip-math-api.md index 37efafbbbf..b3698ff2b3 100644 --- a/projects/hip/docs/markdown/hip-math-api.md +++ b/projects/hip/docs/markdown/hip-math-api.md @@ -1433,7 +1433,7 @@ __device__ float __expf(float x); __device__ static float __fadd_rd(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fadd_rn @@ -1441,7 +1441,7 @@ __device__ static float __fadd_rd(float x, float y); __device__ static float __fadd_rn(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fadd_ru @@ -1449,7 +1449,7 @@ __device__ static float __fadd_rn(float x, float y); __device__ static float __fadd_ru(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fadd_rz @@ -1457,7 +1457,7 @@ __device__ static float __fadd_ru(float x, float y); __device__ static float __fadd_rz(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fdiv_rd @@ -1465,7 +1465,7 @@ __device__ static float __fadd_rz(float x, float y); __device__ static float __fdiv_rd(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fdiv_rn @@ -1473,7 +1473,7 @@ __device__ static float __fdiv_rd(float x, float y); __device__ static float __fdiv_rn(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fdiv_ru @@ -1481,7 +1481,7 @@ __device__ static float __fdiv_rn(float x, float y); __device__ static float __fdiv_ru(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fdiv_rz @@ -1489,7 +1489,7 @@ __device__ static float __fdiv_ru(float x, float y); __device__ static float __fdiv_rz(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fdividef @@ -1505,7 +1505,7 @@ __device__ static float __fdividef(float x, float y); __device__ float __fmaf_rd(float x, float y, float z); ``` -**Description:** Supported +**Description:** Unsupported ### __fmaf_rn @@ -1513,7 +1513,7 @@ __device__ float __fmaf_rd(float x, float y, float z); __device__ float __fmaf_rn(float x, float y, float z); ``` -**Description:** Supported +**Description:** Unsupported ### __fmaf_ru @@ -1521,7 +1521,7 @@ __device__ float __fmaf_rn(float x, float y, float z); __device__ float __fmaf_ru(float x, float y, float z); ``` -**Description:** Supported +**Description:** Unsupported ### __fmaf_rz @@ -1529,7 +1529,7 @@ __device__ float __fmaf_ru(float x, float y, float z); __device__ float __fmaf_rz(float x, float y, float z); ``` -**Description:** Supported +**Description:** Unsupported ### __fmul_rd @@ -1537,7 +1537,7 @@ __device__ float __fmaf_rz(float x, float y, float z); __device__ static float __fmul_rd(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fmul_rn @@ -1545,7 +1545,7 @@ __device__ static float __fmul_rd(float x, float y); __device__ static float __fmul_rn(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fmul_ru @@ -1553,7 +1553,7 @@ __device__ static float __fmul_rn(float x, float y); __device__ static float __fmul_ru(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fmul_rz @@ -1561,7 +1561,7 @@ __device__ static float __fmul_ru(float x, float y); __device__ static float __fmul_rz(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __frcp_rd @@ -1609,7 +1609,7 @@ __device__ float __frsqrt_rn(float x); __device__ float __fsqrt_rd(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __fsqrt_rn @@ -1617,7 +1617,7 @@ __device__ float __fsqrt_rd(float x); __device__ float __fsqrt_rn(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __fsqrt_ru @@ -1625,7 +1625,7 @@ __device__ float __fsqrt_rn(float x); __device__ float __fsqrt_ru(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __fsqrt_rz @@ -1633,7 +1633,7 @@ __device__ float __fsqrt_ru(float x); __device__ float __fsqrt_rz(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __fsub_rd @@ -1641,7 +1641,7 @@ __device__ float __fsqrt_rz(float x); __device__ static float __fsub_rd(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fsub_rn @@ -1649,7 +1649,7 @@ __device__ static float __fsub_rd(float x, float y); __device__ static float __fsub_rn(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported ### __fsub_ru @@ -1657,7 +1657,15 @@ __device__ static float __fsub_rn(float x, float y); __device__ static float __fsub_ru(float x, float y); ``` -**Description:** Supported +**Description:** Unsupported + + +### __fsub_rz +```cpp +__device__ static float __fsub_rz(float x, float y); + +``` +**Description:** Unsupported ### __log10f @@ -1729,7 +1737,7 @@ __device__ float __tanf(float x); __device__ static double __dadd_rd(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dadd_rn @@ -1737,7 +1745,7 @@ __device__ static double __dadd_rd(double x, double y); __device__ static double __dadd_rn(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dadd_ru @@ -1745,7 +1753,7 @@ __device__ static double __dadd_rn(double x, double y); __device__ static double __dadd_ru(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dadd_rz @@ -1753,7 +1761,7 @@ __device__ static double __dadd_ru(double x, double y); __device__ static double __dadd_rz(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __ddiv_rd @@ -1761,7 +1769,7 @@ __device__ static double __dadd_rz(double x, double y); __device__ static double __ddiv_rd(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __ddiv_rn @@ -1769,7 +1777,7 @@ __device__ static double __ddiv_rd(double x, double y); __device__ static double __ddiv_rn(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __ddiv_ru @@ -1777,7 +1785,7 @@ __device__ static double __ddiv_rn(double x, double y); __device__ static double __ddiv_ru(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __ddiv_rz @@ -1785,7 +1793,7 @@ __device__ static double __ddiv_ru(double x, double y); __device__ static double __ddiv_rz(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dmul_rd @@ -1793,7 +1801,7 @@ __device__ static double __ddiv_rz(double x, double y); __device__ static double __dmul_rd(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dmul_rn @@ -1801,7 +1809,7 @@ __device__ static double __dmul_rd(double x, double y); __device__ static double __dmul_rn(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dmul_ru @@ -1809,7 +1817,7 @@ __device__ static double __dmul_rn(double x, double y); __device__ static double __dmul_ru(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dmul_rz @@ -1817,7 +1825,7 @@ __device__ static double __dmul_ru(double x, double y); __device__ static double __dmul_rz(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __drcp_rd @@ -1857,7 +1865,7 @@ __device__ double __drcp_rz(double x); __device__ double __dsqrt_rd(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __dsqrt_rn @@ -1865,7 +1873,7 @@ __device__ double __dsqrt_rd(double x); __device__ double __dsqrt_rn(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __dsqrt_ru @@ -1873,7 +1881,7 @@ __device__ double __dsqrt_rn(double x); __device__ double __dsqrt_ru(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __dsqrt_rz @@ -1881,7 +1889,7 @@ __device__ double __dsqrt_ru(double x); __device__ double __dsqrt_rz(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __dsub_rd @@ -1889,7 +1897,7 @@ __device__ double __dsqrt_rz(double x); __device__ static double __dsub_rd(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dsub_rn @@ -1897,7 +1905,7 @@ __device__ static double __dsub_rd(double x, double y); __device__ static double __dsub_rn(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dsub_ru @@ -1905,7 +1913,7 @@ __device__ static double __dsub_rn(double x, double y); __device__ static double __dsub_ru(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __dsub_rz @@ -1913,7 +1921,7 @@ __device__ static double __dsub_ru(double x, double y); __device__ static double __dsub_rz(double x, double y); ``` -**Description:** Supported +**Description:** Unsupported ### __fma_rd @@ -1921,7 +1929,7 @@ __device__ static double __dsub_rz(double x, double y); __device__ double __fma_rd(double x, double y, double z); ``` -**Description:** Supported +**Description:** Unsupported ### __fma_rn @@ -1929,7 +1937,7 @@ __device__ double __fma_rd(double x, double y, double z); __device__ double __fma_rn(double x, double y, double z); ``` -**Description:** Supported +**Description:** Unsupported ### __fma_ru @@ -1937,7 +1945,7 @@ __device__ double __fma_rn(double x, double y, double z); __device__ double __fma_ru(double x, double y, double z); ``` -**Description:** Supported +**Description:** Unsupported ### __fma_rz @@ -1945,7 +1953,7 @@ __device__ double __fma_ru(double x, double y, double z); __device__ double __fma_rz(double x, double y, double z); ``` -**Description:** Supported +**Description:** Unsupported ### __brev From eb4053e3fecaeb5dd3952dfaafc93957d08038c2 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 6 Nov 2018 16:32:14 +0000 Subject: [PATCH 8/9] Guard the OCML rounded operations instead Instead of commenting all these functions out, guard the functions with a macro OCML_BASIC_ROUNDED_OPERATIONS. [ROCm/hip commit: e59c33250ab9aa4195e14f89aac3779be2f34bcd] --- .../include/hip/hcc_detail/math_functions.h | 346 +++++++++--------- .../hipDoublePrecisionIntrinsics.cpp | 52 +-- .../hipSinglePrecisionIntrinsics.cpp | 54 +-- 3 files changed, 236 insertions(+), 216 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 8a6091858b..08be321d68 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -514,69 +514,73 @@ float __exp10f(float x) { return __ocml_exp10_f32(x); } __DEVICE__ inline float __expf(float x) { return __ocml_exp_f32(x); } -// __DEVICE__ -// inline -// float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); } -// __DEVICE__ -// inline -// float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); } -// __DEVICE__ -// inline -// float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); } -// __DEVICE__ -// inline -// float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } -// __DEVICE__ -// inline -// float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); } -// __DEVICE__ -// inline -// float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); } -// __DEVICE__ -// inline -// float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); } -// __DEVICE__ -// inline -// float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); } +#if defined OCML_BASIC_ROUNDED_OPERATIONS +__DEVICE__ +inline +float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); } +__DEVICE__ +inline +float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); } +__DEVICE__ +inline +float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); } +__DEVICE__ +inline +float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } +__DEVICE__ +inline +float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); } +__DEVICE__ +inline +float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); } +__DEVICE__ +inline +float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); } +__DEVICE__ +inline +float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); } +#endif __DEVICE__ inline float __fdividef(float x, float y) { return x / y; } -// __DEVICE__ -// inline -// float __fmaf_rd(float x, float y, float z) -// { -// return __ocml_fma_rtn_f32(x, y, z); -// } -// __DEVICE__ -// inline -// float __fmaf_rn(float x, float y, float z) -// { -// return __ocml_fma_rte_f32(x, y, z); -// } -// __DEVICE__ -// inline -// float __fmaf_ru(float x, float y, float z) -// { -// return __ocml_fma_rtp_f32(x, y, z); -// } -// __DEVICE__ -// inline -// float __fmaf_rz(float x, float y, float z) -// { -// return __ocml_fma_rtz_f32(x, y, z); -// } -// __DEVICE__ -// inline -// float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); } -// __DEVICE__ -// inline -// float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); } -// __DEVICE__ -// inline -// float __fmul_ru(float x, float y) { return __ocml_mul_rtp_f32(x, y); } -// __DEVICE__ -// inline -// float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); } +#if defined OCML_BASIC_ROUNDED_OPERATIONS +__DEVICE__ +inline +float __fmaf_rd(float x, float y, float z) +{ + return __ocml_fma_rtn_f32(x, y, z); +} +__DEVICE__ +inline +float __fmaf_rn(float x, float y, float z) +{ + return __ocml_fma_rte_f32(x, y, z); +} +__DEVICE__ +inline +float __fmaf_ru(float x, float y, float z) +{ + return __ocml_fma_rtp_f32(x, y, z); +} +__DEVICE__ +inline +float __fmaf_rz(float x, float y, float z) +{ + return __ocml_fma_rtz_f32(x, y, z); +} +__DEVICE__ +inline +float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); } +__DEVICE__ +inline +float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); } +__DEVICE__ +inline +float __fmul_ru(float x, float y) { return __ocml_mul_rtp_f32(x, y); } +__DEVICE__ +inline +float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); } +#endif __DEVICE__ inline float __frcp_rd(float x) { return __llvm_amdgcn_rcp_f32(x); } @@ -592,30 +596,32 @@ float __frcp_rz(float x) { return __llvm_amdgcn_rcp_f32(x); } __DEVICE__ inline float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); } -// __DEVICE__ -// inline -// float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); } -// __DEVICE__ -// inline -// float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); } -// __DEVICE__ -// inline -// float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); } -// __DEVICE__ -// inline -// float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); } -// __DEVICE__ -// inline -// float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); } -// __DEVICE__ -// inline -// float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); } -// __DEVICE__ -// inline -// float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); } -// __DEVICE__ -// inline -// float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); } +#if defined OCML_BASIC_ROUNDED_OPERATIONS +__DEVICE__ +inline +float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); } +__DEVICE__ +inline +float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); } +__DEVICE__ +inline +float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); } +__DEVICE__ +inline +float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); } +__DEVICE__ +inline +float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); } +__DEVICE__ +inline +float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); } +__DEVICE__ +inline +float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); } +__DEVICE__ +inline +float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); } +#endif __DEVICE__ inline float __log10f(float x) { return __ocml_log10_f32(x); } @@ -1034,42 +1040,44 @@ double yn(int n, double x) } // BEGIN INTRINSICS -// __DEVICE__ -// inline -// double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); } -// __DEVICE__ -// inline -// double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); } -// __DEVICE__ -// inline -// double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); } -// __DEVICE__ -// inline -// double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } -// __DEVICE__ -// inline -// double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); } -// __DEVICE__ -// inline -// double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); } -// __DEVICE__ -// inline -// double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); } -// __DEVICE__ -// inline -// double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); } -// __DEVICE__ -// inline -// double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); } -// __DEVICE__ -// inline -// double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); } -// __DEVICE__ -// inline -// double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); } -// __DEVICE__ -// inline -// double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); } +#if defined OCML_BASIC_ROUNDED_OPERATIONS +__DEVICE__ +inline +double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); } +__DEVICE__ +inline +double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); } +__DEVICE__ +inline +double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); } +__DEVICE__ +inline +double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } +__DEVICE__ +inline +double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); } +__DEVICE__ +inline +double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); } +__DEVICE__ +inline +double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); } +__DEVICE__ +inline +double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); } +__DEVICE__ +inline +double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); } +__DEVICE__ +inline +double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); } +__DEVICE__ +inline +double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); } +__DEVICE__ +inline +double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); } +#endif __DEVICE__ inline double __drcp_rd(double x) { return __llvm_amdgcn_rcp_f64(x); } @@ -1082,54 +1090,56 @@ double __drcp_ru(double x) { return __llvm_amdgcn_rcp_f64(x); } __DEVICE__ inline double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); } -// __DEVICE__ -// inline -// double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); } -// __DEVICE__ -// inline -// double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); } -// __DEVICE__ -// inline -// double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); } -// __DEVICE__ -// inline -// double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); } -// __DEVICE__ -// inline -// double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); } -// __DEVICE__ -// inline -// double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); } -// __DEVICE__ -// inline -// double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); } -// __DEVICE__ -// inline -// double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); } -// __DEVICE__ -// inline -// double __fma_rd(double x, double y, double z) -// { -// return __ocml_fma_rtn_f64(x, y, z); -// } -// __DEVICE__ -// inline -// double __fma_rn(double x, double y, double z) -// { -// return __ocml_fma_rte_f64(x, y, z); -// } -// __DEVICE__ -// inline -// double __fma_ru(double x, double y, double z) -// { -// return __ocml_fma_rtp_f64(x, y, z); -// } -// __DEVICE__ -// inline -// double __fma_rz(double x, double y, double z) -// { -// return __ocml_fma_rtz_f64(x, y, z); -// } +#if defined OCML_BASIC_ROUNDED_OPERATIONS +__DEVICE__ +inline +double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); } +__DEVICE__ +inline +double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); } +__DEVICE__ +inline +double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); } +__DEVICE__ +inline +double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); } +__DEVICE__ +inline +double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); } +__DEVICE__ +inline +double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); } +__DEVICE__ +inline +double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); } +__DEVICE__ +inline +double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); } +__DEVICE__ +inline +double __fma_rd(double x, double y, double z) +{ + return __ocml_fma_rtn_f64(x, y, z); +} +__DEVICE__ +inline +double __fma_rn(double x, double y, double z) +{ + return __ocml_fma_rte_f64(x, y, z); +} +__DEVICE__ +inline +double __fma_ru(double x, double y, double z) +{ + return __ocml_fma_rtp_f64(x, y, z); +} +__DEVICE__ +inline +double __fma_rz(double x, double y, double z) +{ + return __ocml_fma_rtz_f64(x, y, z); +} +#endif // END INTRINSICS // END DOUBLE diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index f6c515c03a..295fd83708 100644 --- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -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) { diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index 623ea08a94..db60099558 100644 --- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp @@ -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); From 8bf242966383755b079fe27d3c37abfe95220594 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 6 Nov 2018 19:53:28 +0000 Subject: [PATCH 9/9] Guard rcp rounded implementation as well Since rcp implementations of non-default rounded versions are not correct or supported in OCML, guard them using the same macro OCML_BASIC_ROUNDED_OPERATIONS. Also update the docs and tests. [ROCm/hip commit: 0cfaa52d15575822a62472d30166714094f0cb29] --- projects/hip/docs/markdown/hip-math-api.md | 18 +++++++++--------- .../include/hip/hcc_detail/math_functions.h | 4 ---- .../deviceLib/hipDoublePrecisionIntrinsics.cpp | 2 -- .../hip/tests/src/deviceLib/hipFloatMath.cpp | 10 ++++++---- .../deviceLib/hipSinglePrecisionIntrinsics.cpp | 2 -- 5 files changed, 15 insertions(+), 21 deletions(-) diff --git a/projects/hip/docs/markdown/hip-math-api.md b/projects/hip/docs/markdown/hip-math-api.md index b3698ff2b3..9b8a3f2f11 100644 --- a/projects/hip/docs/markdown/hip-math-api.md +++ b/projects/hip/docs/markdown/hip-math-api.md @@ -1569,7 +1569,7 @@ __device__ static float __fmul_rz(float x, float y); __device__ float __frcp_rd(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __frcp_rn @@ -1577,7 +1577,7 @@ __device__ float __frcp_rd(float x); __device__ float __frcp_rn(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __frcp_ru @@ -1585,7 +1585,7 @@ __device__ float __frcp_rn(float x); __device__ float __frcp_ru(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __frcp_rz @@ -1593,7 +1593,7 @@ __device__ float __frcp_ru(float x); __device__ float __frcp_rz(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __frsqrt_rn @@ -1601,7 +1601,7 @@ __device__ float __frcp_rz(float x); __device__ float __frsqrt_rn(float x); ``` -**Description:** Supported +**Description:** Unsupported ### __fsqrt_rd @@ -1833,7 +1833,7 @@ __device__ static double __dmul_rz(double x, double y); __device__ double __drcp_rd(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __drcp_rn @@ -1841,7 +1841,7 @@ __device__ double __drcp_rd(double x); __device__ double __drcp_rn(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __drcp_ru @@ -1849,7 +1849,7 @@ __device__ double __drcp_rn(double x); __device__ double __drcp_ru(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __drcp_rz @@ -1857,7 +1857,7 @@ __device__ double __drcp_ru(double x); __device__ double __drcp_rz(double x); ``` -**Description:** Supported +**Description:** Unsupported ### __dsqrt_rd diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 08be321d68..557257b2b0 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -580,7 +580,6 @@ float __fmul_ru(float x, float y) { return __ocml_mul_rtp_f32(x, y); } __DEVICE__ inline float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); } -#endif __DEVICE__ inline float __frcp_rd(float x) { return __llvm_amdgcn_rcp_f32(x); } @@ -596,7 +595,6 @@ float __frcp_rz(float x) { return __llvm_amdgcn_rcp_f32(x); } __DEVICE__ inline float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); } -#if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ inline float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); } @@ -1077,7 +1075,6 @@ double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); } __DEVICE__ inline double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); } -#endif __DEVICE__ inline double __drcp_rd(double x) { return __llvm_amdgcn_rcp_f64(x); } @@ -1090,7 +1087,6 @@ double __drcp_ru(double x) { return __llvm_amdgcn_rcp_f64(x); } __DEVICE__ inline double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); } -#if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ inline double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); } diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 295fd83708..939bdae743 100644 --- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -47,12 +47,10 @@ __device__ void double_precision_intrinsics() { __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); -#if defined OCML_BASIC_ROUNDED_OPERATIONS __dsqrt_rd(4.0); __dsqrt_rn(4.0); __dsqrt_ru(4.0); diff --git a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp index ee83309f28..c6a07e26a9 100644 --- a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp +++ b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp @@ -38,11 +38,13 @@ __global__ void floatMath(float* In, float* Out) { Out[tid] = __cosf(In[tid]); Out[tid] = __exp10f(Out[tid]); Out[tid] = __expf(Out[tid]); +#if defined OCML_BASIC_ROUNDED_OPERATIONS Out[tid] = __frsqrt_rn(Out[tid]); - //Out[tid] = __fsqrt_rd(Out[tid]); - //Out[tid] = __fsqrt_rn(Out[tid]); - //Out[tid] = __fsqrt_ru(Out[tid]); - //Out[tid] = __fsqrt_rz(Out[tid]); + Out[tid] = __fsqrt_rd(Out[tid]); + Out[tid] = __fsqrt_rn(Out[tid]); + Out[tid] = __fsqrt_ru(Out[tid]); + Out[tid] = __fsqrt_rz(Out[tid]); +#endif Out[tid] = __log10f(Out[tid]); Out[tid] = __log2f(Out[tid]); Out[tid] = __logf(Out[tid]); diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index db60099558..b216b3cb54 100644 --- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp @@ -59,13 +59,11 @@ __device__ void single_precision_intrinsics() { __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); -#if defined OCML_BASIC_ROUNDED_OPERATIONS __fsqrt_rd(4.0f); __fsqrt_rn(4.0f); __fsqrt_ru(4.0f);