diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index fb453d02e9..6455fa6cd1 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -30,6 +30,7 @@ THE SOFTWARE. #include #include #include +#include typedef unsigned long ulong; typedef unsigned int uint; @@ -970,4 +971,62 @@ unsigned __smid(void) #endif //defined(__clang__) && defined(__HIP__) +// 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, unsigned char 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) { + unsigned char val8 = static_cast(val); + return __hip_hc_memset(ptr, val8, size); +} #endif diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 13eaf4fda4..c8ac3b4ef4 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -260,25 +260,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/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index 34e198c61d..87fbe0fcbc 100644 --- a/hipamd/src/device_util.cpp +++ b/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; -} - // abort __device__ void abort() { return hc::abort(); }