Move __hip_hc_memcpy and __hip_hc_memset from device_utils.cpp to device_functions.h as inline functions

[ROCm/hip commit: 8136a348ab]
2018-06-22 14:22:26 -04:00
parent 4c4de52f5a
commit 244c9875e2
3 changed files with 58 additions and 65 deletions
@@ -92,58 +92,6 @@ __device__ void* __hip_hc_free(void* ptr) {
    return nullptr;
 }

-
-// loop unrolling
-__device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size) {
-    auto dstPtr = static_cast<uint8_t*>(dst);
-    auto srcPtr = static_cast<const uint8_t*>(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;
-}
-
-__device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) {
-    auto dstPtr = static_cast<uint8_t*>(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;
-}
-
 __device__ long long int clock64() { return (long long int)hc::__cycle_u64(); };
 __device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); };