[PATCH] D104904: [OpenMP][AMDGCN] Initial math headers support
Jon Chesterfield via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 25 11:12:03 PDT 2021
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
More information about the cfe-commits
mailing list