diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h index 8474f066df..7c5a2f2e36 100644 --- a/include/hcc_detail/hip_runtime.h +++ b/include/hcc_detail/hip_runtime.h @@ -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) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 4f95320ac3..d4a6857559 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -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 //----