[PATCH] D104904: [OpenMP][AMDGCN] Initial math headers support

[Jon Chesterfield via Phabricator via cfe-commits]

JonChesterfield added inline comments.


================
Comment at: clang/lib/Headers/__clang_hip_cmath.h:29
+#ifdef __OPENMP_AMDGCN__
+#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
+#define __constant__ __attribute__((constant))
----------------
scchan wrote:
> `__DEVICE__` should not imply constexpr.  It should be added to each function separately.
iirc rocm does that with a macro called __DEVICE_NOCE__, perhaps we could go with __DEVICE_CONSTEXPR__. There's some interaction with overloading rules and different glibc versions, so it would be nice to tag exactly the same functions as constexpr on nvptx and amdgcn


================
Comment at: clang/lib/Headers/__clang_hip_math.h:29
+#else
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#endif
----------------
wonder if HIP would benefit from nothrow here


================
Comment at: clang/lib/Headers/__clang_hip_math.h:35
+#ifdef __OPENMP_AMDGCN__
+#define __RETURN_TYPE int
+#else
----------------
I'd expect openmp to match the cplusplus/c distinction here, as openmp works on C source


================
Comment at: clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h:93
+
+#define __device__ __attribute__((device))
+
----------------
i think this should be `#define __device__`


================
Comment at: clang/test/Headers/Inputs/include/cstdlib:15
 
+#ifndef __AMDGCN__
 namespace std
----------------
I think I'd expect builtin_labs et al to work on amdgcn, are we missing lowering for them?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D104904/new/

https://reviews.llvm.org/D104904