From 5d857b2bc3daf9680705d90bece41fcef0beb167 Mon Sep 17 00:00:00 2001
From: streamhsa <sandeep.kumar@amd.com>
Date: Wed, 10 Feb 2016 20:01:16 +0800
Subject: [PATCH] Resolved test issues

[ROCm/clr commit: 2f8d56e90389972fb7668dcd4857737f4bac0e47]
---
 projects/clr/hipamd/tests/src/CMakeLists.txt  |  2 +-
 .../clr/hipamd/tests/src/hipGridLaunch.cpp    |  2 +-
 projects/clr/hipamd/tests/src/hip_anyall.cpp  | 18 ++++++++++------
 projects/clr/hipamd/tests/src/hip_ballot.cpp  | 21 +++++++++++--------
 4 files changed, 26 insertions(+), 17 deletions(-)

diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt
index ba4be66a22..0ec287b334 100644
--- a/projects/clr/hipamd/tests/src/CMakeLists.txt
+++ b/projects/clr/hipamd/tests/src/CMakeLists.txt
@@ -126,7 +126,7 @@ make_test(hipEventRecord --iterations 10)
 make_test(hipMemset " " )
 make_test(hipMemset --N 10    --memsetval 0x42 )  # small copy, just 10 bytes.
 make_test(hipMemset --N 10013 --memsetval 0x5a )  # oddball size.
-make_test(hipMemset --N 500M  --memsetval 0xa6 )  # big copy
+make_test(hipMemset --N 256M  --memsetval 0xa6 )  # big copy
 make_test(hipGridLaunch " " )
 
 make_test(hipMemcpy  " " )
diff --git a/projects/clr/hipamd/tests/src/hipGridLaunch.cpp b/projects/clr/hipamd/tests/src/hipGridLaunch.cpp
index 4502446b3b..f13781362e 100644
--- a/projects/clr/hipamd/tests/src/hipGridLaunch.cpp
+++ b/projects/clr/hipamd/tests/src/hipGridLaunch.cpp
@@ -37,7 +37,7 @@ __device__ int foo(int i)
 //Syntax we would like to support with GRID_LAUNCH enabled:
 template <typename T>
 __global__ void
-vectorADD2( grid_launch_parm lp,
+vectorADD2( hipLaunchParm lp,
             T *A_d,
             T *B_d,
             T *C_d,
diff --git a/projects/clr/hipamd/tests/src/hip_anyall.cpp b/projects/clr/hipamd/tests/src/hip_anyall.cpp
index e126541766..c6ac36e2e3 100644
--- a/projects/clr/hipamd/tests/src/hip_anyall.cpp
+++ b/projects/clr/hipamd/tests/src/hip_anyall.cpp
@@ -39,12 +39,17 @@ __global__ void
 
 
 int main(int argc, char *argv[])
-{
-
+{ int warpSize;
+  hipDeviceProp_t devProp;
+  hipDeviceGetProperties(&devProp, 0);
+  if(strncmp(devProp.name,"Fiji",1)==0)  warpSize =64;
+  else warpSize =32;
+  int anycount =0;
+  int allcount =0;
   int Num_Threads_per_Block      = 1024;
   int Num_Blocks_per_Grid        = 1;
-  int Num_Warps_per_Block        = Num_Threads_per_Block/64;
-  int Num_Warps_per_Grid         = (Num_Threads_per_Block*Num_Blocks_per_Grid)/64;
+  int Num_Warps_per_Block        = Num_Threads_per_Block/warpSize;
+  int Num_Warps_per_Grid         = (Num_Threads_per_Block*Num_Blocks_per_Grid)/warpSize;
   
   int * host_any  = ( int*)malloc(Num_Warps_per_Grid*sizeof(int));
   int * host_all  = ( int*)malloc(Num_Warps_per_Grid*sizeof(int));
@@ -69,10 +74,11 @@ for (int i=0; i<Num_Warps_per_Grid; i++)
 
     printf("warp no. %d __any = %d \n",i,host_any[i]);
     printf("warp no. %d __all = %d \n",i,host_all[i]);
-
+    if (host_any[i]!=1) ++anycount;
+    if (host_all[i]!=1) ++allcount; 
 
 }
-
+if (anycount == 0 && allcount ==1) printf("PASSED"); else printf("FAILED");
 
   return EXIT_SUCCESS;
 
diff --git a/projects/clr/hipamd/tests/src/hip_ballot.cpp b/projects/clr/hipamd/tests/src/hip_ballot.cpp
index 41087ca56c..4693e58875 100644
--- a/projects/clr/hipamd/tests/src/hip_ballot.cpp
+++ b/projects/clr/hipamd/tests/src/hip_ballot.cpp
@@ -15,16 +15,19 @@ __global__ void
 
 
 int main(int argc, char *argv[])
-{
-
+{ int warpSize;
+  hipDeviceProp_t devProp;
+  hipDeviceGetProperties(&devProp, 0);
+  if(strncmp(devProp.name,"Fiji",1)==0)  warpSize =64;
+  else warpSize =32;
   unsigned int Num_Threads_per_Block      = 512;
   unsigned int Num_Blocks_per_Grid        = 1;
-  unsigned int Num_Warps_per_Block        = Num_Threads_per_Block/64;
-  unsigned int Num_Warps_per_Grid         = (Num_Threads_per_Block*Num_Blocks_per_Grid)/64;
+  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; 
   HIP_ASSERT(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid*sizeof(unsigned int)));
-
+  int divergent_count =0;
   for (int i=0; i<Num_Warps_per_Grid; i++) host_ballot[i] = 0;
 
  
@@ -36,13 +39,13 @@ int main(int argc, char *argv[])
   HIP_ASSERT(hipMemcpy(host_ballot, device_ballot, Num_Warps_per_Grid*sizeof(unsigned int), hipMemcpyDeviceToHost));
   for (int i=0; i<Num_Warps_per_Grid; i++) {
 
-     if ((host_ballot[i] == 0)||(host_ballot[i]/64 == 64)) std::cout << "Warp " << i << " IS convergent- Predicate true for " << host_ballot[i]/64 << " threads\n";
-
-     else std::cout << "Warp " << i << " IS divergent - Predicate true for " << host_ballot[i]/64<< " threads\n";
+     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";
+	  divergent_count++;}
 }
 
-
+if (divergent_count==1) printf("PASSED"); else printf("FAILED");
   return EXIT_SUCCESS;
 
 }