diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index ae1b96c979..17d908be6d 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -210,4 +210,62 @@ __device__ char4 __hip_hc_sub8pk(char4, char4); __device__ char4 __hip_hc_mul8pk(char4, char4); +// loop unrolling +static inline __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size) { + auto dstPtr = static_cast(dst); + auto srcPtr = static_cast(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; +} + +static inline __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) { + auto dstPtr = static_cast(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; +} +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) { + uint8_t val8 = static_cast(val); + return __hip_hc_memset(ptr, val8, size); +} #endif diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index f67da8d51e..57192fdf82 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -364,25 +364,12 @@ static constexpr Coordinates threadIdx; #endif // defined __HCC__ #if __HCC_OR_HIP_CLANG__ -extern "C" __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size); -extern "C" __device__ void* __hip_hc_memset(void* ptr, uint8_t val, size_t size); extern "C" __device__ void* __hip_hc_malloc(size_t); extern "C" __device__ void* __hip_hc_free(void* ptr); static inline __device__ void* malloc(size_t size) { return __hip_hc_malloc(size); } - static inline __device__ void* free(void* ptr) { return __hip_hc_free(ptr); } -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) { - uint8_t val8 = static_cast(val); - return __hip_hc_memset(ptr, val8, size); -} - - #ifdef __HCC_ACCELERATOR__ #ifdef HC_FEATURE_PRINTF diff --git a/projects/clr/hipamd/src/device_util.cpp b/projects/clr/hipamd/src/device_util.cpp index 613e35f0cc..5f37e0dab4 100644 --- a/projects/clr/hipamd/src/device_util.cpp +++ b/projects/clr/hipamd/src/device_util.cpp @@ -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(dst); - auto srcPtr = static_cast(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(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(); };