Workaround problem with fusion in CUDA 9

[ptrendx]

Description

Fixes #17020

The problem comes from the bug in how NVRTC in CUDA 9 handles the default-device flag. That flag is supposed to mark all the functions in the file as __device__ functions, but it should leave the functions decorated differently (like kernels decorated with __global__) alone. This is the behavior in CUDA 10+. In CUDA 9, however, this __device__ attribute is applied to every function (including kernels), which is incompatible with __launch_bounds__() attribute that we use for kernels.

This PR removes the usage of default-device flag for NVRTC compilation and instead manually decorates all the required functions as __device__

[DickJC123]

This looks like a straightforward fix to the nvrtc problem. I'm sure it was tedious to add __device__ in the many places required. What might help the file size and code duplication is the macro:

// Create a fast-math function by appending an 'f' to the canonical name.
#define DEFINE_FASTMATH_FUNC(func)  \
  template <typename DType> \
  __device__ inline DType func(const DType val) { \
    return func ## f(val); \
  }

DEFINE_FASTMATH_FUNC(sin)
... etc

This only applies to some of the many functions defined in fused_op-inl.h, so I approve this PR independent of this suggestion.

[perdasilva]

Awesome stuff! Thank you @ptrendx and @DickJC123 - this puts CD back on track ^^