From ca142c6d9cf78570a3661918c60ba31549f89542 Mon Sep 17 00:00:00 2001
From: scchan <siuchi.chan@amd.com>
Date: Mon, 1 Feb 2016 23:55:31 -0600
Subject: [PATCH] adding shfl, shfl_up, shfl_down, shfl_xor intrinsics

[ROCm/hip commit: 04f3e3e59887a0098224513ed3011d689e1dd272]
---
 projects/hip/include/hcc_detail/hip_runtime.h | 41 +++++++++++++++++++
 1 file changed, 41 insertions(+)

diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h
index 688c9e6959..734596a476 100644
--- a/projects/hip/include/hcc_detail/hip_runtime.h
+++ b/projects/hip/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.