cms-patatrack / cmssw

CMSSW fork of the Patatrack project
https://patatrack.web.cern.ch/patatrack/index.html
Apache License 2.0
2 stars 5 forks source link

Validate edm::isFinite() / edm::isNotFinite() on GPU and compare with std::isfinite() #617

Open ericcano opened 3 years ago

ericcano commented 3 years ago

edm::isNotFinite has not yet been validate on GPU. It should be and compared with its std:: counterpart in both CPU and GPU cases.

It seems benign at first sight. It implies a few integer operations (while the std:: version calls a more open ended intrinsic (I did not dig out what the std intrinsic actually contains, in any of CPU of GPU):

  template <typename T>
  constexpr bool isNotFinite(T x) {
    return !isFinite(x);
  }

  template <>
  constexpr bool isFinite(float x) {
    const unsigned int mask = 0x7f800000;
    union {
      unsigned int l;
      float d;
    } v = {.d = x};
    return (v.l & mask) != mask;
  }

While std:: version is:

  template<typename _Tp>
    inline typename __gnu_cxx::__enable_if<__is_arithmetic<_Tp>::__value,
                       int>::__type
    isfinite(_Tp __f)
    {
      typedef typename __gnu_cxx::__promote<_Tp>::__type __type;
      return __builtin_isfinite(__type(__f));
    }
fwyzard commented 3 years ago

What about using CUDA's isfinite() math function ? __device__ ​ __RETURN_TYPE isfinite ( float a ) __device__ ​ __RETURN_TYPE isfinite ( double a )