From 725486fb118c6eb5fad8d30dfd13845fba875561 Mon Sep 17 00:00:00 2001 From: Nico Trost Date: Thu, 21 Mar 2019 10:30:10 +0100 Subject: [PATCH 1/2] fixed loss of accuracy in hipCfma() --- hipamd/include/hip/hcc_detail/hip_complex.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_complex.h b/hipamd/include/hip/hcc_detail/hip_complex.h index d19abd7247..128e2d670b 100644 --- a/hipamd/include/hip/hcc_detail/hip_complex.h +++ b/hipamd/include/hip/hcc_detail/hip_complex.h @@ -328,8 +328,8 @@ __device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r) { - float real = (p.x * q.x) + r.x; - float imag = (q.x * p.y) + r.y; + double real = (p.x * q.x) + r.x; + double imag = (q.x * p.y) + r.y; real = -(p.y * q.y) + real; imag = (p.x * q.y) + imag; From cf7ad0f184ddb3cea07e0ae5725cc06dc4f3f79e Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Thu, 21 Mar 2019 17:18:16 +0000 Subject: [PATCH 2/2] HACK for SWDEV-173477 For code objects with global symbols of length 0, ROCR runtime would ignore them even though they exist in the symbol table. Therefore the result from read_agent_globals() can't be trusted entirely. As a workaround to tame applications which depend on the existence of global symbols with length 0, always return hipSuccess here. This behavior shall be reverted once ROCR runtime has been fixed to address SWDEV-173477 --- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index f10c7ef575..3c6792e884 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -2621,7 +2621,20 @@ hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), it0->second.cend(), name); - return *dptr ? hipSuccess : hipErrorNotFound; + // HACK for SWDEV-173477 + // + // For code objects with global symbols of length 0, ROCR runtime would + // ignore them even though they exist in the symbol table. Therefore the + // result from read_agent_globals() can't be trusted entirely. + // + // As a workaround to tame applications which depend on the existence of + // global symbols with length 0, always return hipSuccess here. + // + // This behavior shall be reverted once ROCR runtime has been fixed to + // address SWDEV-173477 + + //return *dptr ? hipSuccess : hipErrorNotFound; + return hipSuccess; } inline