From 2afa40873a1e8cd5ef779696306e354d4336a76f Mon Sep 17 00:00:00 2001 From: cjatin Date: Tue, 16 Feb 2021 14:21:24 +0530 Subject: [PATCH] SWDEV-272811 - Using OCML functions for conversion of fp to int Change-Id: Id0c02bf6a2e0553d522035b078a9db4cd3fad215 --- include/hip/amd_detail/device_functions.h | 72 +++++++++++++++-------- 1 file changed, 48 insertions(+), 24 deletions(-) diff --git a/include/hip/amd_detail/device_functions.h b/include/hip/amd_detail/device_functions.h index c7cba34313..40b06df812 100644 --- a/include/hip/amd_detail/device_functions.h +++ b/include/hip/amd_detail/device_functions.h @@ -781,32 +781,44 @@ __device__ static inline int __double2loint(double x) { return tmp[0]; } -__device__ static inline int __double2int_rd(double x) { return (int)x; } -__device__ static inline int __double2int_rn(double x) { return (int)x; } -__device__ static inline int __double2int_ru(double x) { return (int)x; } +__device__ static inline int __double2int_rd(double x) { return (int)__ocml_floor_f64(x); } +__device__ static inline int __double2int_rn(double x) { return (int)__ocml_rint_f64(x); } +__device__ static inline int __double2int_ru(double x) { return (int)__ocml_ceil_f64(x); } __device__ static inline int __double2int_rz(double x) { return (int)x; } -__device__ static inline long long int __double2ll_rd(double x) { return (long long int)x; } -__device__ static inline long long int __double2ll_rn(double x) { return (long long int)x; } -__device__ static inline long long int __double2ll_ru(double x) { return (long long int)x; } -__device__ static inline long long int __double2ll_rz(double x) { return (long long int)x; } +__device__ static inline long long int __double2ll_rd(double x) { + return (long long)__ocml_floor_f64(x); +} +__device__ static inline long long int __double2ll_rn(double x) { + return (long long)__ocml_rint_f64(x); +} +__device__ static inline long long int __double2ll_ru(double x) { + return (long long)__ocml_ceil_f64(x); +} +__device__ static inline long long int __double2ll_rz(double x) { return (long long)x; } -__device__ static inline unsigned int __double2uint_rd(double x) { return (unsigned int)x; } -__device__ static inline unsigned int __double2uint_rn(double x) { return (unsigned int)x; } -__device__ static inline unsigned int __double2uint_ru(double x) { return (unsigned int)x; } +__device__ static inline unsigned int __double2uint_rd(double x) { + return (unsigned int)__ocml_floor_f64(x); +} +__device__ static inline unsigned int __double2uint_rn(double x) { + return (unsigned int)__ocml_rint_f64(x); +} +__device__ static inline unsigned int __double2uint_ru(double x) { + return (unsigned int)__ocml_ceil_f64(x); +} __device__ static inline unsigned int __double2uint_rz(double x) { return (unsigned int)x; } __device__ static inline unsigned long long int __double2ull_rd(double x) { - return (unsigned long long int)x; + return (unsigned long long int)__ocml_floor_f64(x); } __device__ static inline unsigned long long int __double2ull_rn(double x) { - return (unsigned long long int)x; + return (unsigned long long int)__ocml_rint_f64(x); } __device__ static inline unsigned long long int __double2ull_ru(double x) { - return (unsigned long long int)x; + return (unsigned long long int)__ocml_ceil_f64(x); } __device__ static inline unsigned long long int __double2ull_rz(double x) { - return (unsigned long long int)x; + return (unsigned long long int)x; } __device__ static inline long long int __double_as_longlong(double x) { @@ -837,27 +849,39 @@ __device__ static inline int __float2int_rn(float x) { return (int)__ocml_rint_f __device__ static inline int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); } __device__ static inline int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); } -__device__ static inline long long int __float2ll_rd(float x) { return (long long int)x; } -__device__ static inline long long int __float2ll_rn(float x) { return (long long int)x; } -__device__ static inline long long int __float2ll_ru(float x) { return (long long int)x; } +__device__ static inline long long int __float2ll_rd(float x) { + return (long long int)__ocml_floor_f32(x); +} +__device__ static inline long long int __float2ll_rn(float x) { + return (long long int)__ocml_rint_f32(x); +} +__device__ static inline long long int __float2ll_ru(float x) { + return (long long int)__ocml_ceil_f32(x); +} __device__ static inline long long int __float2ll_rz(float x) { return (long long int)x; } -__device__ static inline unsigned int __float2uint_rd(float x) { return (unsigned int)x; } -__device__ static inline unsigned int __float2uint_rn(float x) { return (unsigned int)x; } -__device__ static inline unsigned int __float2uint_ru(float x) { return (unsigned int)x; } +__device__ static inline unsigned int __float2uint_rd(float x) { + return (unsigned int)__ocml_floor_f32(x); +} +__device__ static inline unsigned int __float2uint_rn(float x) { + return (unsigned int)__ocml_rint_f32(x); +} +__device__ static inline unsigned int __float2uint_ru(float x) { + return (unsigned int)__ocml_ceil_f32(x); +} __device__ static inline unsigned int __float2uint_rz(float x) { return (unsigned int)x; } __device__ static inline unsigned long long int __float2ull_rd(float x) { - return (unsigned long long int)x; + return (unsigned long long int)__ocml_floor_f32(x); } __device__ static inline unsigned long long int __float2ull_rn(float x) { - return (unsigned long long int)x; + return (unsigned long long int)__ocml_rint_f32(x); } __device__ static inline unsigned long long int __float2ull_ru(float x) { - return (unsigned long long int)x; + return (unsigned long long int)__ocml_ceil_f32(x); } __device__ static inline unsigned long long int __float2ull_rz(float x) { - return (unsigned long long int)x; + return (unsigned long long int)x; } __device__ static inline int __float_as_int(float x) {