diff --git a/hipamd/docs/markdown/hip_faq.md b/hipamd/docs/markdown/hip_faq.md
index f8a7489fb7..101d99486e 100644
--- a/hipamd/docs/markdown/hip_faq.md
+++ b/hipamd/docs/markdown/hip_faq.md
@@ -239,4 +239,7 @@ $HIP/include/hip/hcc_detail/host_defines.h
 Or pass "-DGENERIC_GRID_LAUNCH=0" to hipcc at application compilation time.
 
 ### What is maximum limit of Generic Grid Launch parameters (grid and block)?
-Product of (grid.x and block.x), (grid.y and block.y) or (grid.z and block.z) should always be less than UINT_MAX.
\ No newline at end of file
+Product of (grid.x and block.x), (grid.y and block.y) or (grid.z and block.z) should always be less than UINT_MAX.
+
+### Are __shfl_*_sync functions supported on HIP platform?
+__shfl_*_sync is not supported on HIP but for nvcc path CUDA 9.0 and above all shuffle calls get redirected to it's sync version.
\ No newline at end of file
diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h
index d908f72a04..d13c999351 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h
@@ -1375,10 +1375,12 @@ hipError_t hipHostFree(void* ptr);
  *
  *  For hipMemcpy, the copy is always performed by the current device (set by hipSetDevice).
  *  For multi-gpu or peer-to-peer configurations, it is recommended to set the current device to the
- * device where the src data is physically located. For optimal peer-to-peer copies, the copy device
- * must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy
- * agent as the current device and src/dest as the peerDevice argument.  if this is not done, the
- * hipMemcpy will still work, but will perform the copy using a staging buffer on the host.
+ *  device where the src data is physically located. For optimal peer-to-peer copies, the copy device
+ *  must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy
+ *  agent as the current device and src/dest as the peerDevice argument.  if this is not done, the
+ *  hipMemcpy will still work, but will perform the copy using a staging buffer on the host.
+ *  Calling hipMemcpy with dst and src pointers that do not match the hipMemcpyKind results in
+ *  undefined behavior.
  *
  *  @param[out]  dst Data being copy to
  *  @param[in]  src Data being copy from
@@ -2089,13 +2091,14 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
 /**
  *  @brief Copies data between host and device.
  *
- *  @param[in]   dst    Destination memory address
- *  @param[in]   dpitch Pitch of destination memory
- *  @param[in]   src    Source memory address
- *  @param[in]   spitch Pitch of source memory
- *  @param[in]   width  Width of matrix transfer (columns in bytes)
- *  @param[in]   height Height of matrix transfer (rows)
- *  @param[in]   kind   Type of transfer
+ *  @param[in]   dst     Destination memory address
+ *  @param[in]   wOffset Destination starting X offset
+ *  @param[in]   hOffset Destination starting Y offset
+ *  @param[in]   src     Source memory address
+ *  @param[in]   spitch  Pitch of source memory
+ *  @param[in]   width   Width of matrix transfer (columns in bytes)
+ *  @param[in]   height  Height of matrix transfer (rows)
+ *  @param[in]   kind    Type of transfer
  *  @return      #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
  * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
  *
@@ -2108,13 +2111,12 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
 /**
  *  @brief Copies data between host and device.
  *
- *  @param[in]   dst    Destination memory address
- *  @param[in]   dpitch Pitch of destination memory
- *  @param[in]   src    Source memory address
- *  @param[in]   spitch Pitch of source memory
- *  @param[in]   width  Width of matrix transfer (columns in bytes)
- *  @param[in]   height Height of matrix transfer (rows)
- *  @param[in]   kind   Type of transfer
+ *  @param[in]   dst     Destination memory address
+ *  @param[in]   wOffset Destination starting X offset
+ *  @param[in]   hOffset Destination starting Y offset
+ *  @param[in]   src     Source memory address
+ *  @param[in]   count   size in bytes to copy
+ *  @param[in]   kind    Type of transfer
  *  @return      #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
  * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
  *
diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
index 23b8a0619d..fe72f33d65 100644
--- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
@@ -236,6 +236,13 @@ typedef struct cudaResourceViewDesc hipResourceViewDesc;
 #define HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT
 #define HIP_FUNC_ATTRIBUTE_MAX CU_FUNC_ATTRIBUTE_MAX
 
+#if CUDA_VERSION >= 9000
+#define __shfl(...)      __shfl_sync(0xffffffff, __VA_ARGS__)
+#define __shfl_up(...)   __shfl_up_sync(0xffffffff, __VA_ARGS__)
+#define __shfl_down(...) __shfl_down_sync(0xffffffff, __VA_ARGS__)
+#define __shfl_xor(...)  __shfl_xor_sync(0xffffffff, __VA_ARGS__)
+#endif // CUDA_VERSION >= 9000
+
 inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
     switch (cuError) {
         case cudaSuccess: