adding shfl, shfl_up, shfl_down, shfl_xor intrinsics

[ROCm/clr commit: ed0a4fc1b7]

projects/clr/hipamd/include/hcc_detail/hip_runtime.h

@@ -351,6 +351,47 @@ __device__ inline unsigned long long int __ballot( int input)
 	return hc::__ballot( input);
 }
 
+// warp shuffle functions
+__device__ int __shfl(int input, int lane, int width) 
+{
+  return hc::__shfl(input,lane,width);
+}
+
+__device__ int __shfl_up(int input, unsigned int lane_delta, int width) 
+{
+  return hc::__shfl_up(input,lane_delta,width);
+}
+
+__device__ int __shfl_down(int input, unsigned int lane_delta, int width) 
+{
+  return hc::__shfl_down(input,lane_delta,width);
+}
+
+__device__ int __shfl_xor(int input, int lane_mask, int width) 
+{
+  return hc::__shfl_xor(input,lane_mask,width);
+}
+
+__device__ float __shfl(float input, int lane, int width) 
+{
+  return hc::__shfl(input,lane,width);
+}
+
+__device__ float __shfl_up(float input, unsigned int lane_delta, int width) 
+{
+  return hc::__shfl_up(input,lane_delta,width);
+}
+
+__device__ float __shfl_down(float input, unsigned int lane_delta, int width) 
+{
+  return hc::__shfl_down(input,lane_delta,width);
+}
+
+__device__ float __shfl_xor(float input, int lane_mask, int width) 
+{
+  return hc::__shfl_xor(input,lane_mask,width);
+}
+
 
 #include <hc_math.hpp>
 // TODO: Choose whether default is precise math or fast math based on compilation flag.