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

[Slava Zakharin via flang-commits]

https://github.com/vzakhari updated https://github.com/llvm/llvm-project/pull/85488

>From 5fd58f70a943e73ad9645b89787c514b5f005780 Mon Sep 17 00:00:00 2001
From: Slava Zakharin <szakharin at nvidia.com>
Date: Fri, 15 Mar 2024 16:25:07 -0700
Subject: [PATCH 1/2] [flang][runtime] Get rid of warnings in F18 runtime CUDA
 build.

---
 flang/runtime/freestanding-tools.h | 23 +++++++++++++++++++++++
 flang/runtime/numeric-templates.h  |  7 +++++++
 flang/runtime/tools.cpp            |  2 +-
 flang/runtime/tools.h              |  2 +-
 4 files changed, 32 insertions(+), 2 deletions(-)

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..3d3fbaf70e2d00 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 (Fortran::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..d656e985d6b46d 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));
+      Fortran::runtime::memchr(data, static_cast<int>(ch), chars));
 }
 
 // Copy payload data from one allocated descriptor to another.

>From bac7d00b9028fb662a66371c503e6ed01956e918 Mon Sep 17 00:00:00 2001
From: Slava Zakharin <szakharin at nvidia.com>
Date: Mon, 18 Mar 2024 09:28:16 -0700
Subject: [PATCH 2/2] Addressed review comment.

---
 flang/runtime/tools.cpp | 2 +-
 flang/runtime/tools.h   | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/flang/runtime/tools.cpp b/flang/runtime/tools.cpp
index 3d3fbaf70e2d00..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 (Fortran::runtime::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 d656e985d6b46d..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 *>(
-      Fortran::runtime::memchr(data, static_cast<int>(ch), chars));
+      runtime::memchr(data, static_cast<int>(ch), chars));
 }
 
 // Copy payload data from one allocated descriptor to another.