From 903dfa94eb0da521883359619b149a55a6e34175 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Fri, 22 Jun 2018 14:22:26 -0400 Subject: [PATCH 1/2] Move __hip_hc_memcpy and __hip_hc_memset from device_utils.cpp to device_functions.h as inline functions --- .../include/hip/hcc_detail/device_functions.h | 58 +++++++++++++++++++ hipamd/include/hip/hcc_detail/hip_runtime.h | 13 ----- hipamd/src/device_util.cpp | 52 ----------------- 3 files changed, 58 insertions(+), 65 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index ae1b96c979..17d908be6d 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/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/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index f67da8d51e..57192fdf82 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/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/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index 613e35f0cc..5f37e0dab4 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; -} - __device__ long long int clock64() { return (long long int)hc::__cycle_u64(); }; __device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); }; From 2531815dfec12cc6383ff67d39c1123fd955c126 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 25 Jun 2018 15:55:02 -0400 Subject: [PATCH 2/2] Include stddef.h in device_functions.h and remove use of uint8_t --- hipamd/include/hip/hcc_detail/device_functions.h | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index 17d908be6d..ab14727f11 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -26,6 +26,7 @@ THE SOFTWARE. #include "host_defines.h" #include +#include extern "C" __device__ unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int); extern "C" __device__ signed int __hip_hc_ir_mul24_int(signed int, signed int); @@ -212,8 +213,8 @@ __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); + auto dstPtr = static_cast(dst); + auto srcPtr = static_cast(src); while (size >= 4u) { dstPtr[0] = srcPtr[0]; @@ -237,8 +238,8 @@ static inline __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_ return dst; } -static inline __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) { - auto dstPtr = static_cast(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; @@ -265,7 +266,7 @@ static inline __device__ void* memcpy(void* dst, const void* src, size_t size) { } static inline __device__ void* memset(void* ptr, int val, size_t size) { - uint8_t val8 = static_cast(val); + unsigned char val8 = static_cast(val); return __hip_hc_memset(ptr, val8, size); } #endif