Merge pull request #674 from mangupta/fix_dtests_on_nvcc

[dtests] Fix hipTestClock, hipTestNew, hipTestGlobalVariable, hipSimpleAtomicsTest & hipTestIncludeMath tests on nvcc path

include/hip/nvcc_detail/hip_runtime.h

@@ -118,6 +118,9 @@ typedef int hipLaunchParm;
     }
 #endif
 
+#define __clock() clock()
+#define __clock64() clock64()
+
 #endif
 
 #endif

tests/src/deviceLib/hipSimpleAtomicsTest.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11 --gpu-architecture=sm_60
  * RUN: %t
  * HIT_END
  */
@@ -215,6 +215,7 @@ template<
     typename T, 
     typename enable_if<
         is_same<T, int>{} || is_same<T, unsigned int>{}>::type* = nullptr>
+__device__
 void testKernelSub(T* g_odata) {
     // Atomic subtraction (final should be 0)
     atomicSub(&g_odata[1], 10);
@@ -333,4 +334,4 @@ int main(int argc, char** argv) {
     hipDeviceReset();
     printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!");
     exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
-}
+}

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

tests/src/deviceLib/hipTestIncludeMath.cpp

@@ -30,6 +30,7 @@ THE SOFTWARE.
 // Incorrect implementation causes compilation failure due to conflict
 // declartions.
 
+#include <new>
 #include <hip/math_functions.h>
 
 // Test __HIP_DEVICE_COMPILE__ is defined after math_functions.h
@@ -45,14 +46,6 @@ __device__ __host__ inline void throw_std_bad_alloc()
   #endif
 }
 
-// Test __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ and __HIP_ARCH_HAS_DYNAMIC_PARALLEL__
-// is defined. Eigen HIP/hcc/Half.h __ldg depends on this.
-#if !defined(__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__) ||                            \
-    !defined(__HIP_ARCH_HAS_DYNAMIC_PARALLEL__)
-#error                                                                         \
-    "__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ or __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ not defined"
-#endif
-
 #include <hip/hip_runtime.h>
 #include "test_common.h"
 

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

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