From f48c53534efe766f0cf6565b285ab58a7ce3be2a Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 1 Nov 2016 16:37:33 -0500
Subject: [PATCH] added inter thread data movement intrinsics

Change-Id: I2a8a8ed49429cb7f96439bd28c4b83b5142737df
---
 include/hip/hcc_detail/hip_runtime.h |  9 +++++++++
 src/hip_hcc.cpp                      | 26 ++++++++++++++++++++++++++
 2 files changed, 35 insertions(+)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 1bda07eb7d..65bf42516f 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -560,6 +560,15 @@ extern "C" __device__ void __threadfence(void);
  */
 __device__ void  __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
 
+__device__ unsigned __hip_ds_bpermute(int index, unsigned src);
+__device__ float __hip_ds_bpermutef(int index, float src);
+__device__ unsigned __hip_ds_permute(int index, unsigned src);
+__device__ float __hip_ds_permutef(int index, float src);
+
+__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern);
+__device__ float __hip_ds_swizzlef(float src, int pattern);
+
+__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl);
 
 // doxygen end Fence Fence
 /**
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index c125829739..15fd3a8237 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -180,7 +180,33 @@ __device__ void* __hip_hc_free(void *ptr)
     return nullptr;
 }
 
+__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
+    return hc::__amdgcn_ds_bpermute(index, src);
+}
 
+__device__ float __hip_ds_bpermutef(int index, float src) {
+    return hc::__amdgcn_ds_bpermute(index, src);
+}
+
+__device__ unsigned __hip_ds_permute(int index, unsigned src) {
+    return hc::__amdgcn_ds_permute(index, src);
+}
+
+__device__ float __hip_ds_permutef(int index, float src) {
+    return hc::__amdgcn_ds_permute(index, src);
+}
+
+__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
+    return hc::__amdgcn_ds_swizzle(src, pattern);
+}
+
+__device__ float __hip_ds_swizzlef(float src, int pattern) {
+    return hc::__amdgcn_ds_swizzle(src, pattern);
+}
+
+__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
+    return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
+}
 //=================================================================================================
 // Thread-local storage:
 //=================================================================================================