devicelib: keep <cstring> in device pass; resolve std::mem* via using-decls

The new-hip-tests / HIP7 work needs device-side memset/memcpy plus host
string functions to coexist with libstdc++ <cstring> in single-source HIP.
Two problems had to be solved together:

1. Host build break. <cstring> does 'using ::memset/::memcpy', so *defining*
   std::memset/std::memcpy (as the earlier approach did) collides with it:
   'target of using declaration conflicts with declaration already in scope'.
   That broke host compilation on every lane (pocl, macOS, salami, rusticl),
   and on macOS surfaced as 'call to memcpy is ambiguous' in hip_fp16_gcc.h.
   Fix: instead of redefining them, declare the __device__ memset/memcpy
   overloads at global scope and surface them into std:: with using-decls.
   A duplicate 'using ::memset/::memcpy' is harmless and, placed after the
   __device__ overloads, makes std::memcpy/std::memset resolve to the device
   (__chip_*) versions in device code.

2. Library device pass. Library headers (e.g. rocPRIM) reference strcmp/
   memset/memcpy in __host__ functions (device-property setup), whose bodies
   are still parsed during the device pass. Guarding <cstring> out of device
   compilation (a previous attempt to keep the SPIR-V module small) left those
   identifiers undeclared and broke the device pass -- hipCUB/rocPRIM/rocRAND
   failed to build on the x86 Intel GPU libraries lane. Fix: include <cstring>
   unconditionally. It contributes only declarations, not definitions, so it
   does not enlarge the module; the using-decls above keep device mem* calls
   pointing at __chip_*.

The DG2 IGC kernel-drop that motivated the device-pass <cstring> guard is
tracked/handled separately (PR #1297) and is not reintroduced by this change.

Author: pvelesko

include/hip/devicelib/sync_and_util.hh

@@ -29,13 +29,16 @@
 
 #include <cstddef>
 #include <cstdint>
-// This header declares device-side overloads of memset()/memcpy() at global
-// scope (below). Without the host <cstring> declarations also in scope, host
-// code that calls memset/memcpy/strcmp/etc. (e.g. rocPRIM device headers
-// compiled for the host) sees only the __device__ overload and fails with
-// "call to __device__ function from __host__ function" or "use of undeclared
-// identifier 'strcmp'". Include <cstring> so the host overloads are always
-// visible; HIP host/device overloading keeps the two sets distinct.
+// Include <cstring> unconditionally. Library headers (e.g. rocPRIM) reference
+// strcmp/memset/memcpy/etc. in __host__ functions (e.g. device-property setup),
+// and in single-source HIP those bodies are still parsed during the device
+// pass -- so the host declarations must be visible in BOTH passes. Guarding
+// <cstring> out of device compilation leaves those identifiers undeclared and
+// breaks the device pass. This only brings in declarations, not definitions, so
+// it does not enlarge the SPIR-V module; the __device__ memset/memcpy overloads
+// declared below (surfaced into std:: via using-declarations) ensure device
+// callers still resolve to the __chip_* implementations rather than the host
+// ones.
 #include <cstring>
 
 __device__ constexpr int warpSize = CHIP_DEFAULT_WARP_SIZE;
@@ -148,16 +151,19 @@ extern "C++" inline __device__ void *memcpy(void *dest, const void *src, size_t
   return __chip_memcpy(dest, src, n);
 }
 
-// Expose device-side memset/memcpy in std:: namespace so that std::memset /
-// std::memcpy in device kernels (e.g. rocPRIM) resolve to the __device__
-// overloads rather than the __host__ ones injected by <cstring> above.
+// Expose the device-side memset/memcpy overloads (declared at global scope
+// above) in std:: so that std::memset / std::memcpy in device kernels (e.g.
+// rocPRIM) resolve to the __device__ overloads. We use using-declarations
+// rather than redefining std::memset/std::memcpy: <cstring> (pulled in
+// unconditionally by other host headers such as spdlog, in both host and
+// device passes) already does `using ::memset/::memcpy`, and a *second*
+// definition of std::memset/std::memcpy collides with it ("target of using
+// declaration conflicts with declaration already in scope"). A duplicate
+// using-declaration of the same ::memset/::memcpy is harmless, and because it
+// appears after the __device__ overloads above it also pulls those into std.
 namespace std {
-inline __device__ void *memset(void *ptr, int value, size_t size) {
-  return __chip_memset(ptr, value, size);
-}
-inline __device__ void *memcpy(void *dest, const void *src, size_t n) {
-  return __chip_memcpy(dest, src, n);
-}
+using ::memcpy;
+using ::memset;
 }
 
 extern "C++" inline __device__ unsigned __activemask()