[flang-commits] [flang] f4e90e3 - [flang][runtime] Get rid of warnings in F18 runtime CUDA build. (#85488)

[via flang-commits]

Author: Slava Zakharin
Date: 2024-03-18T16:29:58-07:00
New Revision: f4e90e3f3c02a94a17bd8fe6ad18dbb46f92d027

URL: https://github.com/llvm/llvm-project/commit/f4e90e3f3c02a94a17bd8fe6ad18dbb46f92d027
DIFF: https://github.com/llvm/llvm-project/commit/f4e90e3f3c02a94a17bd8fe6ad18dbb46f92d027.diff

LOG: [flang][runtime] Get rid of warnings in F18 runtime CUDA build. (#85488)

Added: 
    

Modified: 
    flang/runtime/freestanding-tools.h
    flang/runtime/numeric-templates.h
    flang/runtime/tools.cpp
    flang/runtime/tools.h

Removed: 
    


################################################################################
diff  --git a/flang/runtime/freestanding-tools.h b/flang/runtime/freestanding-tools.h
index bdc11ae93ac909..682b4c9b89294b 100644
--- a/flang/runtime/freestanding-tools.h
+++ b/flang/runtime/freestanding-tools.h
@@ -42,6 +42,11 @@
 #define STD_REALLOC_UNSUPPORTED 1
 #endif
 
+#if !defined(STD_MEMCHR_UNSUPPORTED) && \
+    (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__)
+#define STD_MEMCHR_UNSUPPORTED 1
+#endif
+
 namespace Fortran::runtime {
 
 #if STD_FILL_N_UNSUPPORTED
@@ -134,5 +139,23 @@ static inline RT_API_ATTRS void *realloc(void *ptr, std::size_t newByteSize) {
 using std::realloc;
 #endif // !STD_REALLOC_UNSUPPORTED
 
+#if STD_MEMCHR_UNSUPPORTED
+// Provides alternative implementation for std::memchr(), if
+// it is not supported.
+static inline RT_API_ATTRS const void *memchr(
+    const void *ptr, int ch, std::size_t count) {
+  auto buf{reinterpret_cast<const unsigned char *>(ptr)};
+  auto c{static_cast<unsigned char>(ch)};
+  for (; count--; ++buf) {
+    if (*buf == c) {
+      return buf;
+    }
+  }
+  return nullptr;
+}
+#else // !STD_MEMCMP_UNSUPPORTED
+using std::memchr;
+#endif // !STD_MEMCMP_UNSUPPORTED
+
 } // namespace Fortran::runtime
 #endif // FORTRAN_RUNTIME_FREESTANDING_TOOLS_H_

diff  --git a/flang/runtime/numeric-templates.h b/flang/runtime/numeric-templates.h
index ecc3b2654d9652..8ea3daaa57bcf8 100644
--- a/flang/runtime/numeric-templates.h
+++ b/flang/runtime/numeric-templates.h
@@ -122,12 +122,19 @@ template <typename T> struct ABSTy {
   static constexpr RT_API_ATTRS T compute(T x) { return std::abs(x); }
 };
 
+// Suppress the warnings about calling __host__-only
+// 'long double' std::frexp, from __device__ code.
+RT_DIAG_PUSH
+RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN
+
 template <typename T> struct FREXPTy {
   static constexpr RT_API_ATTRS T compute(T x, int *e) {
     return std::frexp(x, e);
   }
 };
 
+RT_DIAG_POP
+
 template <typename T> struct ILOGBTy {
   static constexpr RT_API_ATTRS int compute(T x) { return std::ilogb(x); }
 };

diff  --git a/flang/runtime/tools.cpp b/flang/runtime/tools.cpp
index e653323ed1de03..71022c7a8c179d 100644
--- a/flang/runtime/tools.cpp
+++ b/flang/runtime/tools.cpp
@@ -175,7 +175,7 @@ RT_API_ATTRS void ShallowCopy(const Descriptor &to, const Descriptor &from) {
 
 RT_API_ATTRS char *EnsureNullTerminated(
     char *str, std::size_t length, Terminator &terminator) {
-  if (std::memchr(str, '\0', length) == nullptr) {
+  if (runtime::memchr(str, '\0', length) == nullptr) {
     char *newCmd{(char *)AllocateMemoryOrCrash(terminator, length + 1)};
     std::memcpy(newCmd, str, length);
     newCmd[length] = '\0';

diff  --git a/flang/runtime/tools.h b/flang/runtime/tools.h
index 392e3fc6c89136..df25eb8882335b 100644
--- a/flang/runtime/tools.h
+++ b/flang/runtime/tools.h
@@ -430,7 +430,7 @@ template <>
 inline RT_API_ATTRS const char *FindCharacter(
     const char *data, char ch, std::size_t chars) {
   return reinterpret_cast<const char *>(
-      std::memchr(data, static_cast<int>(ch), chars));
+      runtime::memchr(data, static_cast<int>(ch), chars));
 }
 
 // Copy payload data from one allocated descriptor to another.