Additional typo and extra space fix

hipamd/include/hcc_detail/hip_runtime.h

@@ -283,7 +283,6 @@ __device__ inline unsigned long long int atomicXor(unsigned long long int* addre
 	return (long long int)hc::atomic_fetch_xor((uint64_t*)address,(uint64_t)val);
 }
 
-#ifdef __HCC__
 #include <hc.hpp>
 // integer intrinsic function __poc __clz __ffs __brev
 __device__ inline unsigned int __popc( unsigned int input)
@@ -346,7 +345,8 @@ __device__ inline int __all(  int input)
 
 __device__ inline int __any( int input)
 {
-	return hc::__any( input);
+	if( hc::__any( input)!=0) return 1;
+	else return 0;
 }
 
 
@@ -355,14 +355,53 @@ __device__ inline unsigned long long int __ballot( int input)
 	return hc::__ballot( input);
 }
 
-#endif
+// warp shuffle functions
+__device__ inline int __shfl(int input, int lane, int width)
+{
+  return hc::__shfl(input,lane,width);
+}
+
+__device__ inline int __shfl_up(int input, unsigned int lane_delta, int width)
+{
+  return hc::__shfl_up(input,lane_delta,width);
+}
+
+__device__ inline int __shfl_down(int input, unsigned int lane_delta, int width)
+{
+  return hc::__shfl_down(input,lane_delta,width);
+}
+
+__device__ inline int __shfl_xor(int input, int lane_mask, int width)
+{
+  return hc::__shfl_xor(input,lane_mask,width);
+}
+
+__device__ inline float __shfl(float input, int lane, int width)
+{
+  return hc::__shfl(input,lane,width);
+}
+
+__device__ inline float __shfl_up(float input, unsigned int lane_delta, int width)
+{
+  return hc::__shfl_up(input,lane_delta,width);
+}
+
+__device__ inline float __shfl_down(float input, unsigned int lane_delta, int width)
+{
+  return hc::__shfl_down(input,lane_delta,width);
+}
+
+__device__ inline float __shfl_xor(float input, int lane_mask, int width)
+{
+  return hc::__shfl_xor(input,lane_mask,width);
+}
 
 
-
-#ifdef __HCC_ACCELERATOR__
 #include <hc_math.hpp>
 // TODO: Choose whether default is precise math or fast math based on compilation flag.
+#ifdef __HCC_ACCELERATOR__
 using namespace hc::precise_math;
+#endif
 
 //TODO: Undo this once min/max functions are supported by hc
 inline int min(int arg1, int arg2) __attribute((hc,cpu)) { \
@@ -372,12 +411,24 @@ inline int max(int arg1, int arg2) __attribute((hc,cpu)) { \
 
 
 //TODO - add a couple fast math operations here, the set here will grow :
-__device__ inline float __log2f(float x) {return hc::fast_math::log2(x); };
+__device__ inline float __cosf(float x) {return hc::fast_math::cosf(x); };
+__device__ inline float __expf(float x) {return hc::fast_math::expf(x); };
+__device__ inline float __frsqrt_rn(float x) {return hc::fast_math::rsqrt(x); };
+__device__ inline float __fsqrt_rd(float x) {return hc::fast_math::sqrt(x); };
+__device__ inline float __fsqrt_rn(float x) {return hc::fast_math::sqrt(x); };
+__device__ inline float __fsqrt_ru(float x) {return hc::fast_math::sqrt(x); };
+__device__ inline float __fsqrt_rz(float x) {return hc::fast_math::sqrt(x); };
+__device__ inline float __log10f(float x) {return hc::fast_math::log10f(x); };
+__device__ inline float __log2f(float x) {return hc::fast_math::log2f(x); };
+__device__ inline float __logf(float x) {return hc::fast_math::logf(x); };
 __device__ inline float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); };
-
-#endif
-
-
+__device__ inline void __sincosf(float x, float *s, float *c) {return hc::fast_math::sincosf(x, s, c); };
+__device__ inline float __sinf(float x) {return hc::fast_math::sinf(x); };
+__device__ inline float __tanf(float x) {return hc::fast_math::tanf(x); };
+__device__ inline float __dsqrt_rd(double x) {return hc::fast_math::sqrt(x); };
+__device__ inline float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); };
+__device__ inline float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); };
+__device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
 
 /**
  * Kernel launching