Fix cooperative group tests.

- `gm` should be a global variable instead of auto local variable.

Change-Id: Ibde1e54ee83fa5725f21820b4b792a97bf7df6ad


[ROCm/hip commit: 6bf6480283]

projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp

@@ -36,6 +36,8 @@ THE SOFTWARE.
 
 using namespace cooperative_groups;
 
+static __device__ int gm[2];
+
 static __global__
 void kernel_cg_grid_group_type(int *sizeTestD,
                                int *thdRankTestD,
@@ -55,7 +57,6 @@ void kernel_cg_grid_group_type(int *sizeTestD,
   isValidTestD[gIdx] = gg.is_valid();
 
   // Test sync
-  __device__ int gm[2];
   if (blockIdx.x == 0 && threadIdx.x == 0)
     gm[0] = 10;
   else if (blockIdx.x == 1 && threadIdx.x == 0)

projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp

@@ -36,6 +36,8 @@ THE SOFTWARE.
 
 using namespace cooperative_groups;
 
+static __device__ int gm[2];
+
 static __global__
 void kernel_cg_grid_group_type_via_base_type(int *sizeTestD,
                                              int *thdRankTestD,
@@ -55,7 +57,6 @@ void kernel_cg_grid_group_type_via_base_type(int *sizeTestD,
   isValidTestD[gIdx] = tg.is_valid();
 
   // Test sync
-  __device__ int gm[2];
   if (blockIdx.x == 0 && threadIdx.x == 0)
     gm[0] = 10;
   else if (blockIdx.x == 1 && threadIdx.x == 0)

projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp

@@ -36,6 +36,8 @@ THE SOFTWARE.
 
 using namespace cooperative_groups;
 
+static __device__ int gm[2];
+
 static __global__
 void kernel_cg_grid_group_type_via_public_api(int *sizeTestD,
                                               int *thdRankTestD,
@@ -55,7 +57,6 @@ void kernel_cg_grid_group_type_via_public_api(int *sizeTestD,
   isValidTestD[gIdx] = is_valid(gg);
 
   // Test sync api
-  __device__ int gm[2];
   if (blockIdx.x == 0 && threadIdx.x == 0)
     gm[0] = 10;
   else if (blockIdx.x == 1 && threadIdx.x == 0)