From a7232c7f9e90cd6e76aa039740aac679078bc672 Mon Sep 17 00:00:00 2001
From: streamhsa <sandeep.kumar@amd.com>
Date: Fri, 19 Feb 2016 20:18:03 +0800
Subject: [PATCH] Resolve issues for hip_popc and hip_ballot on nvcc

---
 tests/src/hip_ballot.cpp | 10 +++++++---
 tests/src/hip_popc.cpp   | 23 ++++++++++-------------
 2 files changed, 17 insertions(+), 16 deletions(-)

diff --git a/tests/src/hip_ballot.cpp b/tests/src/hip_ballot.cpp
index 1f3012e689..76fb673ecc 100644
--- a/tests/src/hip_ballot.cpp
+++ b/tests/src/hip_ballot.cpp
@@ -9,7 +9,8 @@ __global__ void
 
    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
    const unsigned int warp_num = hipThreadIdx_x >> pshift;
-   atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popcll(__ballot(tid - 245)));
+   if (pshift ==6) {atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popcll(__ballot(tid - 245)));}
+	else {atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popc(__ballot(tid - 245)));}
  
 }
 
@@ -18,8 +19,11 @@ int main(int argc, char *argv[])
 { int warpSize, pshift;
   hipDeviceProp_t devProp;
   hipDeviceGetProperties(&devProp, 0);
-  if(strncmp(devProp.name,"Fiji",1)==0)  {warpSize =64; pshift =6;}
+  
+  if(strncmp(devProp.name,"Fiji",1)==0)  
+  {warpSize = 64; pshift =6;}
   else {warpSize =32; pshift =5;}
+  
   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;
@@ -41,7 +45,7 @@ int main(int argc, char *argv[])
 
      if ((host_ballot[i] == 0)||(host_ballot[i]/warpSize == warpSize)) std::cout << "Warp " << i << " IS convergent- Predicate true for " << host_ballot[i]/warpSize << " threads\n";
 
-     else {std::cout << "Warp " << i << " IS divergent - Predicate true for " << host_ballot[i]/warpSize<< " threads\n";
+     else {std::cout << " Warp " << i << " IS divergent - Predicate true for " << host_ballot[i]/warpSize<< " threads\n";
 	  divergent_count++;}
 }
 
diff --git a/tests/src/hip_popc.cpp b/tests/src/hip_popc.cpp
index d03c11f32a..0227bdb97c 100644
--- a/tests/src/hip_popc.cpp
+++ b/tests/src/hip_popc.cpp
@@ -31,8 +31,8 @@ THE SOFTWARE.
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
 
-#define WIDTH     32
-#define HEIGHT    32
+#define WIDTH     16
+#define HEIGHT    16
 
 #define NUM       (WIDTH*HEIGHT)
 
@@ -52,10 +52,10 @@ unsigned int popcountCPU( T value) {
 	}
 	return ret;
 }
-
+ 
 __global__ void 
-HIP_kernel(hipLaunchParm lp,
-             unsigned int* a, unsigned int* b,unsigned int* c, unsigned long long int* d, int width, int height) 
+HIP_kernel(hipLaunchParm lp, 
+             unsigned int* a, unsigned int* b,unsigned  int* c, unsigned long long int* d, int width, int height) 
   {
  
       int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
@@ -71,20 +71,18 @@ HIP_kernel(hipLaunchParm lp,
 
 
 
-
-
 using namespace std;
 
 int main() {
   
   unsigned int* hostA;
   unsigned int* hostB;
-  unsigned  int* hostC; 
+  unsigned int* hostC; 
   unsigned long long int* hostD;
 
   unsigned int* deviceA;
   unsigned int* deviceB;
-  unsigned  int* deviceC;
+  unsigned int* deviceC;
   unsigned long long int* deviceD;
 
   hipDeviceProp_t devProp;
@@ -107,7 +105,7 @@ int main() {
   // initialize the input data
   for (i = 0; i < NUM; i++) {
     hostB[i] = i;
-	hostD[i] = 1099511627776-i;
+    hostD[i] = 1099511627776-i;
   }
   
   HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
@@ -128,7 +126,7 @@ int main() {
 
   HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(unsigned int), hipMemcpyDeviceToHost));
   HIP_ASSERT(hipMemcpy(hostC, deviceC, NUM*sizeof(unsigned int), hipMemcpyDeviceToHost));
-  // verify the results
+   // verify the results
   errors = 0;
   for (i = 0; i < NUM; i++) {
 	printf("gpu_popc =%d, cpu_popc =%d \n",hostA[i],popcountCPU(hostB[i]));
@@ -149,11 +147,10 @@ int main() {
     }
   }
   if (errors!=0) {
-    printf("FAILED: %d errors\n",errors);
+      printf("FAILED: %d errors\n",errors);
   } else {
       printf ("__popcll() PASSED!\n");
   }
-
   HIP_ASSERT(hipFree(deviceA));
   HIP_ASSERT(hipFree(deviceB));
   HIP_ASSERT(hipFree(deviceC));