[clr] rely on __builtins for memset/memcpy device functions (#4047)

## Motivation

Simplify the device memset/memcpy code

## Technical Details

Use compiler builtins for memset and memcpy. Compiler has recently made
some changes to make sure that builtins generate good code for these
operations which made it to upstream LLVM, so we can use this from now.
This also changes the signature of memset function, from unsigned char
to something that builtin expects, this should not result in any ABI
break since `unsigned char` basically fits in an `int` (standard int
promotion) and since its a `__device__` function it will be inlined and
will have no traces as a function inside final code.

## JIRA ID

NA

## Test Plan


## Test Result


## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Author: cjatin

projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h

@@ -888,65 +888,23 @@ unsigned __smid(void)
 
 #endif  // defined(__clang__) && defined(__HIP__)
 
-
-// loop unrolling
+// rely on `__builtin_* functions for memcpy/memset
 static inline __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size) {
-  auto dstPtr = static_cast<unsigned char*>(dst);
-  auto srcPtr = static_cast<const unsigned char*>(src);
-
-  while (size >= 4u) {
-    dstPtr[0] = srcPtr[0];
-    dstPtr[1] = srcPtr[1];
-    dstPtr[2] = srcPtr[2];
-    dstPtr[3] = srcPtr[3];
-
-    size -= 4u;
-    srcPtr += 4u;
-    dstPtr += 4u;
-  }
-  switch (size) {
-    case 3:
-      dstPtr[2] = srcPtr[2];
-    case 2:
-      dstPtr[1] = srcPtr[1];
-    case 1:
-      dstPtr[0] = srcPtr[0];
-  }
-
-  return dst;
+  return __builtin_memcpy(dst, src, size);
 }
 
-static inline __device__ void* __hip_hc_memset(void* dst, unsigned char val, size_t size) {
-  auto dstPtr = static_cast<unsigned char*>(dst);
-
-  while (size >= 4u) {
-    dstPtr[0] = val;
-    dstPtr[1] = val;
-    dstPtr[2] = val;
-    dstPtr[3] = val;
-
-    size -= 4u;
-    dstPtr += 4u;
-  }
-  switch (size) {
-    case 3:
-      dstPtr[2] = val;
-    case 2:
-      dstPtr[1] = val;
-    case 1:
-      dstPtr[0] = val;
-  }
-
-  return dst;
+// change the value from unsigned char to int, what a builtin expects.
+static inline __device__ void* __hip_hc_memset(void* dst, int val, size_t size) {
+  return __builtin_memset(dst, val, size);
 }
+
 #ifndef __OPENMP_AMDGCN__
 static inline __device__ void* memcpy(void* dst, const void* src, size_t size) {
   return __hip_hc_memcpy(dst, src, size);
 }
 
 static inline __device__ void* memset(void* ptr, int val, size_t size) {
-  unsigned char val8 = static_cast<unsigned char>(val);
-  return __hip_hc_memset(ptr, val8, size);
+  return __hip_hc_memset(ptr, val, size);
 }
 #endif  // !__OPENMP_AMDGCN__