From 0bae7eaad33a5ccadbbbb11e85167362d648f955 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Mon, 17 Sep 2018 15:32:05 +0530
Subject: [PATCH] [dests] Fix hipTestClock, hipTestNew & hipTestGlobalVariable
 tests for nvcc

nvcc does not support global kernels in struct/class

Change-Id: I2d7297e0c3725564215e20dbdd31c0bb8d7a07de


[ROCm/clr commit: f63ffaf6e5c91dcbe9301cc4ae6e32f33c30eaf5]
---
 .../tests/src/deviceLib/hipTestClock.cpp       |  5 +----
 .../hipamd/tests/src/deviceLib/hipTestNew.cpp  |  4 +---
 .../tests/src/kernel/hipTestGlobalVariable.cpp | 18 ++++++------------
 3 files changed, 8 insertions(+), 19 deletions(-)

diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp
index 46f64e35a3..ee6dca8a42 100644
--- a/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp
+++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp
@@ -33,8 +33,6 @@ THE SOFTWARE.
 #define LEN 512
 #define SIZE 2048
 
-struct TestClock {
-     
   static __global__ void kernel1(int* Ad) {
       int tid = threadIdx.x + blockIdx.x * blockDim.x;
       Ad[tid] = clock() + clock64() + __clock() + __clock64();
@@ -61,9 +59,8 @@ struct TestClock {
         assert(0 != A[i]);
     }
   }
-};
 
 int main() {
-  TestClock().run();
+  run();
   passed();
 }
diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp
index 60774ff21d..d644f8b483 100644
--- a/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp
+++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp
@@ -33,7 +33,6 @@ THE SOFTWARE.
 #define LEN 512
 #define SIZE 2048
 
-struct TestPlacementNew {
   class A {
   public:
       __device__ A() {
@@ -63,9 +62,8 @@ struct TestPlacementNew {
         assert(i == A[i]);
     }
   }
-};
 
 int main() {
-  TestPlacementNew().run();
+  run();
   passed();
 }
diff --git a/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp b/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp
index 2209e2c254..8ab8bef9c2 100644
--- a/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp
+++ b/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp
@@ -33,15 +33,14 @@ THE SOFTWARE.
 #define LEN 512
 #define SIZE 2048
 
-struct TestConstantGlobalVar {
-  static __constant__ int ConstantGlobalVar;
+  __constant__ int ConstantGlobalVar = 123;
 
   static __global__ void kernel(int* Ad) {
       int tid = threadIdx.x + blockIdx.x * blockDim.x;
       Ad[tid] = ConstantGlobalVar;
   }
 
-  void run() {
+  void runTestConstantGlobalVar() {
     int *A, *Ad;
     A = new int[LEN];
     for (unsigned i = 0; i < LEN; i++) {
@@ -56,11 +55,8 @@ struct TestConstantGlobalVar {
         assert(123 == A[i]);
     }
   }
-};
-__constant__ int TestConstantGlobalVar::ConstantGlobalVar = 123;
 
-struct TestGlobalArray {
-  static __device__ int GlobalArray[LEN];
+  __device__ int GlobalArray[LEN];
 
   static __global__ void kernelWrite() {
       int tid = threadIdx.x + blockIdx.x * blockDim.x;
@@ -71,7 +67,7 @@ struct TestGlobalArray {
       Ad[tid] = GlobalArray[tid];
   }
 
-  void run() {
+  void runTestGlobalArray() {
     int *A, *Ad;
     A = new int[LEN];
     for (unsigned i = 0; i < LEN; i++) {
@@ -87,11 +83,9 @@ struct TestGlobalArray {
         assert(i == A[i]);
     }
   }
-};
-__device__ int TestGlobalArray::GlobalArray[LEN];
 
 int main() {
-  TestConstantGlobalVar().run();
-  TestGlobalArray().run();
+  runTestConstantGlobalVar();
+  runTestGlobalArray();
   passed();
 }