Add support for few CUDA compatible reduction operations like __syncthreads_count()

Change-Id: I0917820058f0edad5c0bdd42dea37f06353fc63c

hipamd/include/hip/hcc_detail/device_functions.h

@@ -1131,6 +1131,30 @@ void __syncthreads()
   __barrier(__CLK_LOCAL_MEM_FENCE);
 }
 
+__device__
+inline
+__attribute__((convergent))
+int __syncthreads_count(int predicate)
+{
+  return __ockl_wgred_add_i32(!!predicate);
+}
+
+__device__
+inline
+__attribute__((convergent))
+int __syncthreads_and(int predicate)
+{
+  return __ockl_wgred_and_i32(!!predicate);
+}
+
+__device__
+inline
+__attribute__((convergent))
+int __syncthreads_or(int predicate)
+{
+  return __ockl_wgred_or_i32(!!predicate);
+}
+
 // hip.amdgcn.bc - device routine
 /*
    HW_ID Register bit structure

hipamd/include/hip/hcc_detail/device_library_decls.h

@@ -74,6 +74,11 @@ extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(vo
 
 extern "C" __device__ void __ockl_atomic_add_noret_f32(float*, float);
 
+extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_add_i32(int a);
+extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_and_i32(int a);
+extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_or_i32(int a);
+
+
 // Introduce local address space
 #define __local __attribute__((address_space(3)))