From 7bb5a14a4e347155f26527df537b80d70f46a833 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() --- include/hip/hcc_detail/hip_complex.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index d19abd7247..128e2d670b 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/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 d941f19399e9b4f040aec5d41123e7073913cbc3 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 --- include/hip/hcc_detail/hip_runtime_api.h | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index f10c7ef575..3c6792e884 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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