SWDEV-1 - Stop using ocml rounding functions (#228)

Directly use the builtins. Use the elementwise versions since there's
no implied errno, regardless of -f[no]-math-errno.

I didn't change the cases unnecessarily casting. The bfloat and vector
cases should work directly.

hipamd/include/hip/amd_detail/amd_device_functions.h

@@ -404,41 +404,47 @@ __device__ static inline int __double2loint(double x) {
     return tmp[0];
 }
 
-__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_rd(double x) {
+  return (int)__builtin_elementwise_floor(x);
+}
+__device__ static inline int __double2int_rn(double x) {
+  return (int)__builtin_elementwise_rint(x);
+}
+__device__ static inline int __double2int_ru(double x) {
+  return (int)__builtin_elementwise_ceil(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)__ocml_floor_f64(x);
+  return (long long)__builtin_elementwise_floor(x);
 }
 __device__ static inline long long int __double2ll_rn(double x) {
-  return (long long)__ocml_rint_f64(x);
+  return (long long)__builtin_elementwise_rint(x);
 }
 __device__ static inline long long int __double2ll_ru(double x) {
-  return (long long)__ocml_ceil_f64(x);
+  return (long long)__builtin_elementwise_ceil(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)__ocml_floor_f64(x);
+  return (unsigned int)__builtin_elementwise_floor(x);
 }
 __device__ static inline unsigned int __double2uint_rn(double x) {
-  return (unsigned int)__ocml_rint_f64(x);
+  return (unsigned int)__builtin_elementwise_rint(x);
 }
 __device__ static inline unsigned int __double2uint_ru(double x) {
-  return (unsigned int)__ocml_ceil_f64(x);
+  return (unsigned int)__builtin_elementwise_ceil(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)__ocml_floor_f64(x);
+  return (unsigned long long int)__builtin_elementwise_floor(x);
 }
 __device__ static inline unsigned long long int __double2ull_rn(double x) {
-  return (unsigned long long int)__ocml_rint_f64(x);
+  return (unsigned long long int)__builtin_elementwise_rint(x);
 }
 __device__ static inline unsigned long long int __double2ull_ru(double x) {
-  return (unsigned long long int)__ocml_ceil_f64(x);
+  return (unsigned long long int)__builtin_elementwise_ceil(x);
 }
 __device__ static inline unsigned long long int __double2ull_rz(double x) {
   return (unsigned long long int)x;
@@ -466,41 +472,41 @@ CUDA implements half as unsigned short whereas, HIP doesn't.
 
 */
 
-__device__ static inline int __float2int_rd(float x) { return (int)__ocml_floor_f32(x); }
-__device__ static inline int __float2int_rn(float x) { return (int)__ocml_rint_f32(x); }
-__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 int __float2int_rd(float x) { return (int)__builtin_elementwise_floor(x); }
+__device__ static inline int __float2int_rn(float x) { return (int)__builtin_elementwise_rint(x); }
+__device__ static inline int __float2int_ru(float x) { return (int)__builtin_elementwise_ceil(x); }
+__device__ static inline int __float2int_rz(float x) { return (int)__builtin_elementwise_trunc(x); }
 
 __device__ static inline long long int __float2ll_rd(float x) {
-  return (long long int)__ocml_floor_f32(x);
+  return (long long int)__builtin_elementwise_floor(x);
 }
 __device__ static inline long long int __float2ll_rn(float x) {
-  return (long long int)__ocml_rint_f32(x);
+  return (long long int)__builtin_elementwise_rint(x);
 }
 __device__ static inline long long int __float2ll_ru(float x) {
-  return (long long int)__ocml_ceil_f32(x);
+  return (long long int)__builtin_elementwise_ceil(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)__ocml_floor_f32(x);
+  return (unsigned int)__builtin_elementwise_floor(x);
 }
 __device__ static inline unsigned int __float2uint_rn(float x) {
-  return (unsigned int)__ocml_rint_f32(x);
+  return (unsigned int)__builtin_elementwise_rint(x);
 }
 __device__ static inline unsigned int __float2uint_ru(float x) {
-  return (unsigned int)__ocml_ceil_f32(x);
+  return (unsigned int)__builtin_elementwise_ceil(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)__ocml_floor_f32(x);
+  return (unsigned long long int)__builtin_elementwise_floor(x);
 }
 __device__ static inline unsigned long long int __float2ull_rn(float x) {
-  return (unsigned long long int)__ocml_rint_f32(x);
+  return (unsigned long long int)__builtin_elementwise_rint(x);
 }
 __device__ static inline unsigned long long int __float2ull_ru(float x) {
-  return (unsigned long long int)__ocml_ceil_f32(x);
+  return (unsigned long long int)__builtin_elementwise_ceil(x);
 }
 __device__ static inline unsigned long long int __float2ull_rz(float x) {
   return (unsigned long long int)x;

hipamd/include/hip/amd_detail/amd_hip_bf16.h

@@ -1548,7 +1548,7 @@ __BF16_HOST_DEVICE_STATIC__ bool operator>=(const __hip_bfloat162& l, const __hi
  * \brief Calculate ceil of bfloat16
  */
 __BF16_DEVICE_STATIC__ __hip_bfloat16 hceil(const __hip_bfloat16 h) {
-  return __float2bfloat16(__ocml_ceil_f32(__bfloat162float(h)));
+  return __float2bfloat16(__builtin_elementwise_ceil(__bfloat162float(h)));
 }
 
 /**
@@ -1588,7 +1588,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hexp2(const __hip_bfloat16 h) {
  * \brief Calculate floor of bfloat16
  */
 __BF16_DEVICE_STATIC__ __hip_bfloat16 hfloor(const __hip_bfloat16 h) {
-  return __float2bfloat16(__ocml_floor_f32(__bfloat162float(h)));
+  return __float2bfloat16(__builtin_elementwise_floor(__bfloat162float(h)));
 }
 
 /**
@@ -1628,7 +1628,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hrcp(const __hip_bfloat16 h) {
  * \brief Round to nearest int
  */
 __BF16_DEVICE_STATIC__ __hip_bfloat16 hrint(const __hip_bfloat16 h) {
-  return __float2bfloat16(__ocml_rint_f32(__bfloat162float(h)));
+  return __float2bfloat16(__builtin_elementwise_rint(__bfloat162float(h)));
 }
 
 /**
@@ -1660,7 +1660,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat16 hsqrt(const __hip_bfloat16 h) {
  * \brief Calculate truncate of bfloat16
  */
 __BF16_DEVICE_STATIC__ __hip_bfloat16 htrunc(const __hip_bfloat16 h) {
-  return __float2bfloat16(__ocml_trunc_f32(__bfloat162float(h)));
+  return __float2bfloat16(__builtin_elementwise_trunc(__bfloat162float(h)));
 }
 
 /**

hipamd/include/hip/amd_detail/amd_hip_fp16.h

@@ -1593,29 +1593,25 @@ THE SOFTWARE.
             __device__
             __half htrunc(__half x)
             {
-                return __half_raw{
-                    __ocml_trunc_f16(static_cast<__half_raw>(x).data)};
+              return __half_raw{__builtin_elementwise_trunc(static_cast<__half_raw>(x).data)};
             }
             inline
             __device__
             __half hceil(__half x)
             {
-                return __half_raw{
-                    __ocml_ceil_f16(static_cast<__half_raw>(x).data)};
+              return __half_raw{__builtin_elementwise_ceil(static_cast<__half_raw>(x).data)};
             }
             inline
             __device__
             __half hfloor(__half x)
             {
-                return __half_raw{
-                   __ocml_floor_f16(static_cast<__half_raw>(x).data)};
+              return __half_raw{__builtin_elementwise_floor(static_cast<__half_raw>(x).data)};
             }
             inline
             __device__
             __half hrint(__half x)
             {
-                return __half_raw{
-                    __ocml_rint_f16(static_cast<__half_raw>(x).data)};
+              return __half_raw{__builtin_elementwise_rint(static_cast<__half_raw>(x).data)};
             }
             inline
             __device__
@@ -1713,25 +1709,25 @@ THE SOFTWARE.
             __device__
             __half2 h2trunc(__half2 x)
             {
-                return __half2{__ocml_trunc_2f16(x)};
+              return __half2{__builtin_elementwise_trunc(static_cast<__half2_raw>(x).data)};
             }
             inline
             __device__
             __half2 h2ceil(__half2 x)
             {
-                return __half2{__ocml_ceil_2f16(x)};
+              return __half2{__builtin_elementwise_ceil(static_cast<__half2_raw>(x).data)};
             }
             inline
             __device__
             __half2 h2floor(__half2 x)
             {
-                return __half2{__ocml_floor_2f16(x)};
+              return __half2{__builtin_elementwise_floor(static_cast<__half2_raw>(x).data)};
             }
             inline
             __device__
             __half2 h2rint(__half2 x)
             {
-                return __half2{__ocml_rint_2f16(x)};
+              return __half2{__builtin_elementwise_rint(static_cast<__half2_raw>(x).data)};
             }
             inline
             __device__

hipamd/include/hip/amd_detail/device_library_decls.h

@@ -61,11 +61,6 @@ extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort);
 extern "C" __device__ __attribute__((const)) uint __ockl_clz_u32(uint);
 extern "C" __device__ __attribute__((const)) uint64_t __ockl_clz_u64(uint64_t);
 
-extern "C" __device__ __attribute__((const)) float __ocml_floor_f32(float);
-extern "C" __device__ __attribute__((const)) float __ocml_rint_f32(float);
-extern "C" __device__ __attribute__((const)) float __ocml_ceil_f32(float);
-extern "C" __device__ __attribute__((const)) float __ocml_trunc_f32(float);
-
 extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
 extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);