[clang] [CUDA/HIP] Fix errors for device function used in host global initializers (PR #222338)

[via cfe-commits]

ronlieb wrote:

hi, seeing this error:
clang -x hip --cuda-host-only --offload-arch=gfx90a -fsyntax-only -nogpulib ~/tmp/foo.hpp
/home/rlieberm/tmp/foo.hpp:6:1: error: unknown type name '__device__'
    6 | __device__ constexpr unsigned strideCounts() { return 4; }
      | ^
1 error generated when compiling for host.

test case:
// Reduced rocWMMA pattern from io_bearer_base.hpp:
//   static constexpr uint32_t IterationCount = reduce_mult(MatrixLayout::strideCounts());
// Host compile should error after 4735d351182f:
//   reference to __device__ function 'strideCounts' in global initializer
 
__device__ constexpr unsigned strideCounts() { return 4; }
constexpr unsigned reduce_mult(unsigned x) { return x; }
 
struct IOBearerBase {
  static constexpr unsigned IterationCount = reduce_mult(strideCounts());
};


https://github.com/llvm/llvm-project/pull/222338