From e5d029b8a40fe9ef30e48681f79bc1e09e4c11db Mon Sep 17 00:00:00 2001
From: Nick Curtis <arghdos@users.noreply.github.com>
Date: Wed, 26 Feb 2020 04:23:56 -0600
Subject: [PATCH] fix long shuffle implementations for windows (#1895)

Fixes for SWDEV-223694

[ROCm/clr commit: a094c87038d6e34cca7c726647f5825033cb6765]
---
 .../include/hip/hcc_detail/device_functions.h | 20 +++++++++++++++++++
 1 file changed, 20 insertions(+)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
index 46ed53ff87..7096841da8 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
@@ -323,6 +323,7 @@ __device__
 inline
 long __shfl(long var, int src_lane, int width = warpSize)
 {
+    #ifndef _MSC_VER
     static_assert(sizeof(long) == 2 * sizeof(int), "");
     static_assert(sizeof(long) == sizeof(uint64_t), "");
 
@@ -333,6 +334,10 @@ long __shfl(long var, int src_lane, int width = warpSize)
     uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
     long tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
     return tmp1;
+    #else
+    static_assert(sizeof(long) == sizeof(int), "");
+    return static_cast<long>(__shfl(static_cast<int>(var), src_lane, width));
+    #endif
 }
 __device__
 inline
@@ -390,6 +395,7 @@ __device__
 inline
 long __shfl_up(long var, unsigned int lane_delta, int width = warpSize)
 {
+    #ifndef _MSC_VER
     static_assert(sizeof(long) == 2 * sizeof(int), "");
     static_assert(sizeof(long) == sizeof(uint64_t), "");
 
@@ -400,6 +406,10 @@ long __shfl_up(long var, unsigned int lane_delta, int width = warpSize)
     uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
     long tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
     return tmp1;
+    #else
+    static_assert(sizeof(long) == sizeof(int), "");
+    return static_cast<long>(__shfl_up(static_cast<int>(var), lane_delta, width));
+    #endif
 }
 __device__
 inline
@@ -455,6 +465,7 @@ __device__
 inline
 long __shfl_down(long var, unsigned int lane_delta, int width = warpSize)
 {
+    #ifndef _MSC_VER
     static_assert(sizeof(long) == 2 * sizeof(int), "");
     static_assert(sizeof(long) == sizeof(uint64_t), "");
 
@@ -465,6 +476,10 @@ long __shfl_down(long var, unsigned int lane_delta, int width = warpSize)
     uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
     long tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
     return tmp1;
+    #else
+    static_assert(sizeof(long) == sizeof(int), "");
+    return static_cast<long>(__shfl_down(static_cast<int>(var), lane_delta, width));
+    #endif
 }
 __device__
 inline
@@ -520,6 +535,7 @@ __device__
 inline
 long __shfl_xor(long var, int lane_mask, int width = warpSize)
 {
+    #ifndef _MSC_VER
     static_assert(sizeof(long) == 2 * sizeof(int), "");
     static_assert(sizeof(long) == sizeof(uint64_t), "");
 
@@ -530,6 +546,10 @@ long __shfl_xor(long var, int lane_mask, int width = warpSize)
     uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
     long tmp1;  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
     return tmp1;
+    #else
+    static_assert(sizeof(long) == sizeof(int), "");
+    return static_cast<long>(__shfl_down(static_cast<int>(var), lane_delta, width));
+    #endif
 }
 __device__
 inline