AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Move warpSize to header, have shuffles use default warpsize.

Browse Source
This commit is contained in:
Ben Sander
2016-02-15 05:41:09 -06:00
parent 57274850f9
commit db3a63360b
2 changed files with 14 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hcc_detail/hip_runtime.h
+14 -9
View File
@@ -108,6 +108,12 @@ THE SOFTWARE.
#define __HCC_C__
#endif
// TODO - hipify-clang - change to use the function call.
//#define warpSize hc::__wavesize()
const int warpSize = 64;
#define clock_t long long int
__device__ inline long long int clock64() { return (long long int)hc::__clock_u64(); };
__device__ inline clock_t clock() { return (clock_t)hc::__clock_u64(); };
@@ -344,42 +350,42 @@ __device__ inline unsigned long long int __ballot( int input)
}
// warp shuffle functions
__device__ inline int __shfl(int input, int lane, int width)
__device__ inline int __shfl(int input, int lane, int width=warpSize)
{
return hc::__shfl(input,lane,width);
}
__device__ inline int __shfl_up(int input, unsigned int lane_delta, int width)
__device__ inline int __shfl_up(int input, unsigned int lane_delta, int width=warpSize)
{
return hc::__shfl_up(input,lane_delta,width);
}
__device__ inline int __shfl_down(int input, unsigned int lane_delta, int width)
__device__ inline int __shfl_down(int input, unsigned int lane_delta, int width=warpSize)
{
return hc::__shfl_down(input,lane_delta,width);
}
__device__ inline int __shfl_xor(int input, int lane_mask, int width)
__device__ inline int __shfl_xor(int input, int lane_mask, int width=warpSize)
{
return hc::__shfl_xor(input,lane_mask,width);
}
__device__ inline float __shfl(float input, int lane, int width)
__device__ inline float __shfl(float input, int lane, int width=warpSize)
{
return hc::__shfl(input,lane,width);
}
__device__ inline float __shfl_up(float input, unsigned int lane_delta, int width)
__device__ inline float __shfl_up(float input, unsigned int lane_delta, int width=warpSize)
{
return hc::__shfl_up(input,lane_delta,width);
}
__device__ inline float __shfl_down(float input, unsigned int lane_delta, int width)
__device__ inline float __shfl_down(float input, unsigned int lane_delta, int width=warpSize)
{
return hc::__shfl_down(input,lane_delta,width);
}
__device__ inline float __shfl_xor(float input, int lane_mask, int width)
__device__ inline float __shfl_xor(float input, int lane_mask, int width=warpSize)
{
return hc::__shfl_xor(input,lane_mask,width);
}
@@ -438,7 +444,6 @@ __device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
#define hipGridDim_z (hc_get_num_groups(0))
extern int warpSize ;
#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
src/hip_hcc.cpp
-6
View File
@@ -49,8 +49,6 @@ THE SOFTWARE.
//---
// Environment variables:
// TODO-HCC - map this to the HC instruction that uses HSAIL to get the wave size.
int warpSize = 64;
// Intended to distinguish whether an environment variable should be visible only in debug mode, or in debug+release.
//static const int debug = 0;
@@ -169,7 +167,6 @@ public:
void init(unsigned device_index, hc::accelerator acc);
hipError_t getProperties(hipDeviceProp_t* prop);
// TODO- create a copy constructor.
~ihipDevice_t();
};
@@ -213,8 +210,6 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator acc)
this->reset();
};
#if 1
// TODO-remove #ifdef
ihipDevice_t::~ihipDevice_t()
{
if (_null_stream) {
@@ -229,7 +224,6 @@ ihipDevice_t::~ihipDevice_t()
}
hsa_signal_destroy(_copy_signal);
}
#endif
//----