Resolved test issues

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  " " )

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,

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;
 

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;
 
 }