diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index 8eb9d6a46c..bed206c0ed 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -44,21 +44,21 @@ __device__ static inline int __mul24(int x, int y) return __hip_hc_ir_mul24_int(x, y); } __device__ long long int __mul64hi(long long int x, long long int y); -__device__ int __mulhi(int x, int y) +__device__ static inline int __mulhi(int x, int y) { return __hip_hc_ir_mulhi_int(x, y); } __device__ unsigned int __popc( unsigned int x); __device__ unsigned int __popcll( unsigned long long int x); -__device__ int __rhadd(int x, int y) +__device__ static inline int __rhadd(int x, int y) { return (x + y + 1) >> 1; } //__device__ unsigned int __sad(int x, int y, int z); /* -Implemented signed version of sad +Implement signed version of sad */ -__device__ unsigned int __uhadd(unsigned int x, unsigned int y) +__device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) { return (x + y) >> 1; } @@ -71,11 +71,6 @@ __device__ unsigned int __umulhi(unsigned int x, unsigned int y); __device__ unsigned int __urhadd(unsigned int x, unsigned int y); __device__ unsigned int __usad(unsigned int x, unsigned int y, unsigned int z); -// warp vote function __all __any __ballot -__device__ int __all( int input); -__device__ int __any( int input); -__device__ unsigned long long int __ballot( int input); - /* Rounding modes are not yet supported in HIP */ diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index f6967c3445..e911d17ebb 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -420,6 +420,11 @@ __device__ unsigned int atomicInc(unsigned int* address, __device__ unsigned int atomicDec(unsigned int* address, unsigned int val); + // warp vote function __all __any __ballot +__device__ int __all( int input); +__device__ int __any( int input); +__device__ unsigned long long int __ballot( int input); + // warp shuffle functions #ifdef __cplusplus __device__ int __shfl(int input, int lane, int width=warpSize); diff --git a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp index abe0b66e90..f9328e7a2c 100644 --- a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp @@ -20,6 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #include "hip/hip_runtime.h" +#include "hip/device_functions.h" #include "test_common.h" #pragma GCC diagnostic ignored "-Wall" diff --git a/hipamd/tests/src/deviceLib/hip_ballot.cpp b/hipamd/tests/src/deviceLib/hip_ballot.cpp index 236ceb57fe..629e676bc7 100644 --- a/hipamd/tests/src/deviceLib/hip_ballot.cpp +++ b/hipamd/tests/src/deviceLib/hip_ballot.cpp @@ -26,9 +26,11 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" +#include "hip/device_functions.h" + #define HIP_ASSERT(x) (assert((x)==hipSuccess)) -__global__ void +__global__ void gpu_ballot(hipLaunchParm lp, unsigned int* device_ballot, int Num_Warps_per_Block,int pshift) { @@ -39,7 +41,7 @@ __global__ void #else atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popc(__ballot(tid - 245))); #endif - + } @@ -47,24 +49,24 @@ int main(int argc, char *argv[]) { int warpSize, pshift; hipDeviceProp_t devProp; hipGetDeviceProperties(&devProp, 0); - + warpSize = devProp.warpSize; int w = warpSize; - pshift = 0; + pshift = 0; while (w >>= 1) ++pshift; - + unsigned int Num_Threads_per_Block = 512; unsigned int Num_Blocks_per_Grid = 1; unsigned int Num_Warps_per_Block = Num_Threads_per_Block/warpSize; unsigned int Num_Warps_per_Grid = (Num_Threads_per_Block*Num_Blocks_per_Grid)/warpSize; unsigned int* host_ballot = (unsigned int*)malloc(Num_Warps_per_Grid*sizeof(unsigned int)); - unsigned int* device_ballot; + unsigned int* device_ballot; HIP_ASSERT(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid*sizeof(unsigned int))); int divergent_count =0; for (int i=0; i #include #include "hip/hip_runtime.h" - +#include "hip/device_functions.h" #define HIP_ASSERT(x) (assert((x)==hipSuccess)) @@ -181,4 +181,3 @@ int main() { return errors; } - diff --git a/hipamd/tests/src/deviceLib/hip_clz.cpp b/hipamd/tests/src/deviceLib/hip_clz.cpp index 5c60b29a2d..869f4406f5 100644 --- a/hipamd/tests/src/deviceLib/hip_clz.cpp +++ b/hipamd/tests/src/deviceLib/hip_clz.cpp @@ -32,6 +32,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime.h" +#include "hip/device_functions.h" #define HIP_ASSERT(x) (assert((x)==hipSuccess)) #define WIDTH 8 @@ -188,4 +189,3 @@ int main() { return errors; } - diff --git a/hipamd/tests/src/deviceLib/hip_ffs.cpp b/hipamd/tests/src/deviceLib/hip_ffs.cpp index dfdc439a21..ba9bd7b9a0 100644 --- a/hipamd/tests/src/deviceLib/hip_ffs.cpp +++ b/hipamd/tests/src/deviceLib/hip_ffs.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime.h" - +#include "hip/device_functions.h" #define HIP_ASSERT(x) (assert((x)==hipSuccess)) @@ -184,4 +184,3 @@ int main() { return errors; } - diff --git a/hipamd/tests/src/deviceLib/hip_popc.cpp b/hipamd/tests/src/deviceLib/hip_popc.cpp index b40bbd2000..6fe214c7fa 100644 --- a/hipamd/tests/src/deviceLib/hip_popc.cpp +++ b/hipamd/tests/src/deviceLib/hip_popc.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime.h" - +#include "hip/device_functions.h" #define HIP_ASSERT(x) (assert((x)==hipSuccess)) @@ -172,4 +172,3 @@ int main() { return errors; } -