SWDEV-345571 - Revert changes for warp shuffle functions

This reverts commit 609598cd15 and 149afa9444.

Change-Id: I548f15055417453a7fd323a481ed410598830521


[ROCm/clr commit: 6c53a670c1]

projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h

@@ -23,20 +23,6 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_WARP_FUNCTIONS_H
 #define HIP_INCLUDE_HIP_AMD_DETAIL_WARP_FUNCTIONS_H
 
-#if !defined(__HIPCC_RTC__)
-  #if !defined(HIP_NO_HALF)
-    #define HIP_NO_HALF
-  #elif
-    #define ALREADY_DEFINED
-  #endif
-
-  #include "hip/hip_fp16.h"
-
-  #if !defined(ALREADY_DEFINED)
-    #undef HIP_NO_HALF
-  #endif
-#endif
-
 static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE;
 
 __device__
@@ -62,20 +48,6 @@ float __shfl(float var, int src_lane, int width = warpSize) {
 }
 __device__
 inline
-__half __shfl(__half var, int src_lane, int width = warpSize) {
-    union { int i; __half h; } tmp; tmp.h = var;
-    tmp.i = __shfl(tmp.i, src_lane, width);
-    return tmp.h;
-}
-__device__
-inline
-__half2 __shfl(__half2 var, int src_lane, int width = warpSize) {
-    union { int i; __half2 h; } tmp; tmp.h = var;
-    tmp.i = __shfl(tmp.i, src_lane, width);
-    return tmp.h;
-}
-__device__
-inline
 double __shfl(double var, int src_lane, int width = warpSize) {
     static_assert(sizeof(double) == 2 * sizeof(int), "");
     static_assert(sizeof(double) == sizeof(uint64_t), "");
@@ -181,20 +153,6 @@ float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) {
 }
 __device__
 inline
-__half __shfl_up(__half var, unsigned int lane_delta, int width = warpSize) {
-    union { int i; __half h; } tmp; tmp.h = var;
-    tmp.i = __shfl_up(tmp.i, lane_delta, width);
-    return tmp.h;
-}
-__device__
-inline
-__half2 __shfl_up(__half2 var, unsigned int lane_delta, int width = warpSize) {
-    union { int i; __half2 h; } tmp; tmp.h = var;
-    tmp.i = __shfl_up(tmp.i, lane_delta, width);
-    return tmp.h;
-}
-__device__
-inline
 double __shfl_up(double var, unsigned int lane_delta, int width = warpSize) {
     static_assert(sizeof(double) == 2 * sizeof(int), "");
     static_assert(sizeof(double) == sizeof(uint64_t), "");
@@ -301,20 +259,6 @@ float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) {
 }
 __device__
 inline
-__half __shfl_down(__half var, unsigned int lane_delta, int width = warpSize) {
-    union { int i; __half h; } tmp; tmp.h = var;
-    tmp.i = __shfl_down(tmp.i, lane_delta, width);
-    return tmp.h;
-}
-__device__
-inline
-__half2 __shfl_down(__half2 var, unsigned int lane_delta, int width = warpSize) {
-    union { int i; __half2 h; } tmp; tmp.h = var;
-    tmp.i = __shfl_down(tmp.i, lane_delta, width);
-    return tmp.h;
-}
-__device__
-inline
 double __shfl_down(double var, unsigned int lane_delta, int width = warpSize) {
     static_assert(sizeof(double) == 2 * sizeof(int), "");
     static_assert(sizeof(double) == sizeof(uint64_t), "");
@@ -418,20 +362,6 @@ float __shfl_xor(float var, int lane_mask, int width = warpSize) {
 }
 __device__
 inline
-__half __shfl_xor(__half var,  int lane_mask, int width = warpSize) {
-    union { int i; __half h; } tmp; tmp.h = var;
-    tmp.i = __shfl_xor(tmp.i, lane_mask, width);
-    return tmp.h;
-}
-__device__
-inline
-__half2 __shfl_xor(__half2 var,  int lane_mask, int width = warpSize) {
-    union { int i; __half2 h; } tmp; tmp.h = var;
-    tmp.i = __shfl_xor(tmp.i, lane_mask, width);
-    return tmp.h;
-}
-__device__
-inline
 double __shfl_xor(double var, int lane_mask, int width = warpSize) {
     static_assert(sizeof(double) == 2 * sizeof(int), "");
     static_assert(sizeof(double) == sizeof(uint64_t), "");