From 28e1fb0cf70f3b442ff43fa32dba40ab85053088 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Mon, 16 Nov 2020 12:35:20 -0500 Subject: [PATCH] Fix cooperative group tests. - `gm` should be a global variable instead of auto local variable. Change-Id: Ibde1e54ee83fa5725f21820b4b792a97bf7df6ad --- .../src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp | 3 ++- .../cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp | 3 ++- .../cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp | 3 ++- 3 files changed, 6 insertions(+), 3 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp index 79f1cb1c38..facfca655e 100755 --- a/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp +++ b/hipamd/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) diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp index 7407f266dd..7c81cbe300 100755 --- a/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp +++ b/hipamd/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) diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp index cb9d8d7c53..be8b34e83b 100755 --- a/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp +++ b/hipamd/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)