Skip to content

Commit

Permalink
adding shfl, shfl_up, shfl_down, shfl_xor intrinsics
Browse files Browse the repository at this point in the history
  • Loading branch information
scchan committed Feb 2, 2016
1 parent 861cba6 commit 04f3e3e
Showing 1 changed file with 41 additions and 0 deletions.
41 changes: 41 additions & 0 deletions include/hcc_detail/hip_runtime.h
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down

0 comments on commit 04f3e3e

Please sign in to comment.