Merge 'master' into 'amd-master'

Change-Id: Ie456ccae78b50086939e1eaf65a5f768d40b4aca
This commit is contained in:
Jenkins
2019-03-25 04:09:54 -05:00
2 changed files with 16 additions and 3 deletions
+2 -2
View File
@@ -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;
@@ -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