[PATCH] D141555: [CUDA] added cmath wrappers to unbreak CUDA compilation after D79555

[Artem Belevich via Phabricator via cfe-commits]

tra created this revision.
Herald added subscribers: mattd, bixia, yaxunl.
Herald added a project: All.
tra updated this revision to Diff 488733.
tra added a comment.
tra updated this revision to Diff 488734.
tra published this revision for review.
tra added reviewers: jlebar, philnik.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.

Simplified the wrapper.


tra added a comment.

nit fix.



================
Comment at: clang/lib/Headers/cuda_wrappers/cmath:108
+#endif // !__has_constexpr_builtin(__builtin_logb)
+  return ::logb(__x);
+}
----------------
`__builtin_logb()` results in a libcall to logb() which we do not have on the GPU side, so we have to use CUDA-provided function.


================
Comment at: clang/lib/Headers/cuda_wrappers/cmath:153
+#endif // __has_constexpr_builtin(__builtin_scalbln)
+  return ::scalbn(__x, __exp);
+}
----------------
Ditto for __builtin_scalbn()


libc++ introduced a handful of internal functions that may or may not be
constexpr, depending on C++ version. For pre-constexpr variants we must declare
__device__ counterparts. Otherwise the code fails to compile on the GPU side.
See https://reviews.llvm.org/D79555


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D141555

Files:
  clang/lib/Headers/CMakeLists.txt
  clang/lib/Headers/cuda_wrappers/cmath

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D141555.488734.patch
Type: text/x-patch
Size: 4142 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230112/d7193c0b/attachment.bin>