Re-sync with upstream.

tests/src/deviceLib/hipDeviceMalloc.cpp

@@ -0,0 +1,188 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+/* HIT_START
+ * BUILD: %t %s NVCC_OPTIONS -std=c++11
+ * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
+ * HIT_END
+ */
+#include "test_common.h"
+#include <iostream>
+#include <complex>
+
+// Tolerance for error
+const double tolerance = 1e-6;
+const bool verbose = false;
+
+#define BLKDIM_X 64
+#define BLKDIM_Y 1
+#define BLKDIM_Z 1
+#define NUM_BLK_X 1
+#define NUM_BLK_Y 1
+#define NUM_BLK_Z 1
+
+#define LEN (BLKDIM_X * BLKDIM_Y * BLKDIM_Z * NUM_BLK_X * NUM_BLK_Y * NUM_BLK_Z)
+
+#define ALL_FUN \
+  OP(add) \
+  OP(sub) \
+  OP(mul) \
+  OP(div)
+
+#define OP(x) CK_##x,
+enum CalcKind {
+  ALL_FUN
+};
+#undef OP
+
+#define OP(x) case CK_##x: return #x;
+std::string getName(enum CalcKind CK) {
+  switch(CK){
+  ALL_FUN
+  }
+}
+#undef OP
+
+// Calculates function.
+// If the function has one argument, B is ignored.
+#define ONE_ARG(func)                                                                              \
+    case CK_##func:                                                                                \
+        return std::func(A);
+
+template <typename FloatT>
+__device__ __host__ FloatT calc(FloatT A, FloatT B, enum CalcKind CK) {
+    switch (CK) {
+        case CK_add:
+            return A + B;
+        case CK_sub:
+            return A - B;
+        case CK_mul:
+            return A * B;
+        case CK_div:
+            return A / B;
+    }
+}
+
+// Allocate memory in kernel and save the address to pA and pB.
+// Copy value from A, B to allocated memory.
+template <typename FloatT>
+__global__ void kernel_alloc(FloatT* A, FloatT* B, FloatT** pA, FloatT** pB) {
+    int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
+        + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
+        + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
+        * hipBlockDim_y;
+    if (tx == 0) {
+        *pA = (FloatT*)malloc(sizeof(FloatT) * LEN);
+        *pB = (FloatT*)malloc(sizeof(FloatT) * LEN);
+        for (int i = 0; i < LEN; i++) {
+            (*pA)[i] = A[i];
+            (*pB)[i] = B[i];
+      }
+    }
+}
+
+// Do calculation using values saved in allocated memmory. pA, pB are buffers
+// containing the address of the device-side allocated array.
+template <typename FloatT>
+__global__ void kernel_free(FloatT** pA, FloatT** pB, FloatT* C, enum CalcKind CK) {
+    int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
+        + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
+        + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
+        * hipBlockDim_y;
+    C[tx] = calc<FloatT>((*pA)[tx], (*pB)[tx], CK);
+    if (tx == 0) {
+        free(*pA);
+        free(*pB);
+    }
+}
+
+template<typename FloatT>
+void test() {
+    FloatT *A, *Ad, *B, *Bd, *C, *Cd, *D;
+    A = new FloatT[LEN];
+    B = new FloatT[LEN];
+    C = new FloatT[LEN];
+    D = new FloatT[LEN];
+    hipMalloc((void**)&Ad, sizeof(FloatT) * LEN);
+    hipMalloc((void**)&Bd, sizeof(FloatT) * LEN);
+    hipMalloc((void**)&Cd, sizeof(FloatT) * LEN);
+
+    for (uint32_t i = 0; i < LEN; i++) {
+        A[i] = (i + 1) * 1.0f;
+        B[i] = A[i];
+        C[i] = A[i];
+    }
+    hipMemcpy(Ad, A, sizeof(FloatT) * LEN, hipMemcpyHostToDevice);
+    hipMemcpy(Bd, B, sizeof(FloatT) * LEN, hipMemcpyHostToDevice);
+
+    // Run kernel for a calculation kind and verify by comparing with host
+    // calculation result. Returns false if fails.
+    auto test_fun = [&](enum CalcKind CK) {
+      // kernel_alloc allocates memory on device side and initialize it.
+      // kernel_free uses allocated memory from kernel_alloc and does the
+      // calculation then free the memory.
+      // pA and pB are buffers to pass the device-side allocated memory address
+      // from kernel_alloc to kernel_free.
+      FloatT **pA, **pB;
+      hipMalloc((FloatT***)&pA, sizeof(FloatT*));
+      hipMalloc((FloatT***)&pB, sizeof(FloatT*));
+      dim3 blkDim(BLKDIM_X, BLKDIM_Y, BLKDIM_Z);
+      dim3 numBlk(NUM_BLK_X, NUM_BLK_Y, NUM_BLK_Z);
+      hipLaunchKernelGGL(kernel_alloc<FloatT>, numBlk, blkDim, 0, 0,
+          Ad, Bd, pA, pB);
+      hipDeviceSynchronize();
+      hipLaunchKernelGGL(kernel_free<FloatT>, numBlk, blkDim, 0, 0,
+          pA, pB, Cd, CK);
+      hipMemcpy(C, Cd, sizeof(FloatT) * LEN, hipMemcpyDeviceToHost);
+      hipFree(pA);
+      hipFree(pB);
+      for (int i = 0; i < LEN; i++) {
+          FloatT Expected = calc(A[i], B[i], CK);
+          FloatT error = std::abs(C[i] - Expected);
+          if (std::abs(Expected) > tolerance) error /= std::abs(Expected);
+          bool pass = error < tolerance;
+          if (verbose || !pass) {
+              std::cout << "Function: " << getName(CK) << " Operands: " << A[i] << " " << B[i]
+                        << " Result: " << C[i] << " Expected: " << Expected << " Error: " << error
+                        << " Pass: " << pass << std::endl;
+        }
+        if (!pass)
+          return false;
+      }
+      return true;
+    };
+
+#define OP(x) assert(test_fun(CK_##x));
+    ALL_FUN
+#undef OP
+
+    hipFree(Ad);
+    hipFree(Bd);
+    hipFree(Cd);
+    delete[] A;
+    delete[] B;
+    delete[] C;
+    delete[] D;
+}
+
+int main() {
+  test<float>();
+  test<double>();
+  passed();
+  return 0;
+}

tests/src/deviceLib/hipFloatMath.cpp

@@ -40,10 +40,10 @@ __global__ void floatMath(hipLaunchParm lp, float* In, float* Out) {
     Out[tid] = __exp10f(Out[tid]);
     Out[tid] = __expf(Out[tid]);
     Out[tid] = __frsqrt_rn(Out[tid]);
-    Out[tid] = __fsqrt_rd(Out[tid]);
-    Out[tid] = __fsqrt_rn(Out[tid]);
-    Out[tid] = __fsqrt_ru(Out[tid]);
-    Out[tid] = __fsqrt_rz(Out[tid]);
+    //Out[tid] = __fsqrt_rd(Out[tid]);
+    //Out[tid] = __fsqrt_rn(Out[tid]);
+    //Out[tid] = __fsqrt_ru(Out[tid]);
+    //Out[tid] = __fsqrt_rz(Out[tid]);
     Out[tid] = __log10f(Out[tid]);
     Out[tid] = __log2f(Out[tid]);
     Out[tid] = __logf(Out[tid]);

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
  */
@@ -334,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/hipTestFMA.cpp

@@ -0,0 +1,140 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+#include <hip/math_functions.h>
+#include <iostream>
+
+#define HIP_ASSERT(status) assert(status == hipSuccess)
+
+#define LEN 50
+#define SIZE (LEN * sizeof(bool))
+
+  __global__ void kernelTestFMA(bool *Ad) {
+    float f = 1.0f / 3.0f;
+    double d = f;
+    int i = 0;
+    auto Check = [&](bool Cond) { Ad[i++] = Cond; };
+    // f * f + 3.0f will be different if promoted to double.
+    float floatResult = fma(f, f, 3.0f);
+    double doubleResult = fma(d, d, 3.0);
+    Check(floatResult != doubleResult);
+
+    // check promote to float.
+    Check(fma(f, f, 3) == floatResult);
+    Check(fma(f, f, (char)3) == floatResult);
+    Check(fma(f, f, (unsigned char)3) == floatResult);
+    Check(fma(f, f, (short)3) == floatResult);
+    Check(fma(f, f, (unsigned short)3) == floatResult);
+    Check(fma(f, f, (int)3) == floatResult);
+    Check(fma(f, f, (unsigned int)3) == floatResult);
+    Check(fma(f, f, (long)3) == floatResult);
+    Check(fma(f, f, (unsigned long)3) == floatResult);
+    Check(fma(f, f, true) == fma(f, f, 1.0f));
+
+    // check promote to double.
+    Check(fma(d, (double)f, 3) == doubleResult);
+    Check(fma(d, (double)f, (char)3) == doubleResult);
+    Check(fma(d, (double)f, (unsigned char)3) == doubleResult);
+    Check(fma(d, (double)f, (short)3) == doubleResult);
+    Check(fma(d, (double)f, (unsigned short)3) == doubleResult);
+    Check(fma(d, (double)f, (int)3) == doubleResult); 
+    Check(fma(d, (double)f, (unsigned int)3) == doubleResult);
+    Check(fma(d, (double)f, (long)3) == doubleResult);
+    Check(fma(d, (double)f, (unsigned long)3) == doubleResult);
+    Check(fma(d, (double)f, true) == fma((double)f, (double)f, 1.0));
+
+    while (i < LEN)
+      Check(true);
+  }
+
+  void runTestFMA() {
+    bool *Ad;
+    bool A[LEN];
+    for (unsigned i = 0; i < LEN; i++) {
+      A[i] = 0;
+    }
+
+    HIP_ASSERT(hipMalloc((void **)&Ad, SIZE));
+    hipLaunchKernelGGL(kernelTestFMA, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad);
+    HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
+
+    for (unsigned i = 0; i < LEN; i++) {
+      assert(A[i]);
+    }
+  }
+
+    __global__ void kernelTestHalfFMA(bool *Ad) {
+      _Float16 h = (_Float16)(1.0f/3.0f);
+      float f = h;
+      double d = f;
+      int i = 0;
+      auto Check = [&](bool Cond) { Ad[i++] = Cond; };
+      // h * h + 3 will be different if promoted to float.
+      _Float16 halfResult = fma(h, h, (_Float16)3);
+      float floatResult = fma(f, f, 3.0f);
+      double doubleResult = fma(d, d, 3.0);
+      Check(halfResult != floatResult);
+      Check(halfResult != doubleResult);
+
+      // check promote to half.
+      Check(fma(h, h, 3) == halfResult);
+      Check(fma(h, h, (char)3) == halfResult);
+      Check(fma(h, h, (unsigned char)3) == halfResult);
+      Check(fma(h, h, (short)3) == halfResult);
+      Check(fma(h, h, (unsigned short)3) == halfResult);
+      Check(fma(h, h, (int)3) == halfResult);
+      Check(fma(h, h, (unsigned int)3) == halfResult);
+      Check(fma(h, h, (long)3) == halfResult);
+      Check(fma(h, h, (unsigned long)3) == halfResult);
+      Check(fma(h, h, true) == fma(h, h, (_Float16)1));
+
+      while (i < LEN)
+        Check(true);
+    }
+
+  void runTestHalfFMA() {
+    bool *Ad;
+    bool A[LEN];
+    for (unsigned i = 0; i < LEN; i++) {
+      A[i] = 0;
+    }
+
+    HIP_ASSERT(hipMalloc((void **)&Ad, SIZE));
+    hipLaunchKernelGGL(kernelTestHalfFMA, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad);
+    HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
+
+    for (unsigned i = 0; i < LEN; i++) {
+      assert(A[i]);
+    }
+  }
+
+int main() {
+  runTestFMA();
+  runTestHalfFMA();
+  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/deviceLib/hip_trig.cpp

@@ -34,21 +34,26 @@ THE SOFTWARE.
 
 #include <iostream>
 
+#define HIP_ASSERT(x) (assert((x) == hipSuccess))
+
 #define LEN 512
 #define SIZE LEN << 2
 
+#define TEST_DEBUG (0)
+
 __global__ void kernel_trig(hipLaunchParm lp, float* In, float* sin_d, float* cos_d, float* tan_d,
                             float* sin_pd, float* cos_pd) {
     int tid = threadIdx.x + blockIdx.x * blockDim.x;
-    sin_d[tid] = __sinf(In[tid]);
-    cos_d[tid] = __cosf(In[tid]);
-    tan_d[tid] = __tanf(In[tid]);
-    __sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]);
+    sin_d[tid] = sinf(In[tid]);
+    cos_d[tid] = cosf(In[tid]);
+    tan_d[tid] = tanf(In[tid]);
+    sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]);
 }
 
 int main() {
     float *In, *sin_h, *cos_h, *tan_h, *sin_ph, *cos_ph;
     float *In_d, *sin_d, *cos_d, *tan_d, *sin_pd, *cos_pd;
+    int errors = 0;
     In = new float[LEN];
     sin_h = new float[LEN];
     cos_h = new float[LEN];
@@ -63,14 +68,16 @@ int main() {
         sin_ph[i] = 0.0f;
         cos_ph[i] = 0.0f;
     }
-    hipMalloc((void**)&In_d, SIZE);
-    hipMalloc((void**)&sin_d, SIZE);
-    hipMalloc((void**)&cos_d, SIZE);
-    hipMalloc((void**)&tan_d, SIZE);
-    hipMalloc((void**)&sin_pd, SIZE);
-    hipMalloc((void**)&cos_pd, SIZE);
+    HIP_ASSERT(hipMalloc((void**)&In_d, SIZE));
+    HIP_ASSERT(hipMalloc((void**)&sin_d, SIZE));
+    HIP_ASSERT(hipMalloc((void**)&cos_d, SIZE));
+    HIP_ASSERT(hipMalloc((void**)&tan_d, SIZE));
+    HIP_ASSERT(hipMalloc((void**)&sin_pd, SIZE));
+    HIP_ASSERT(hipMalloc((void**)&cos_pd, SIZE));
+
     hipMemcpy(In_d, In, SIZE, hipMemcpyHostToDevice);
-    hipLaunchKernel(kernel_trig, dim3(LEN, 1, 1), dim3(1, 1, 1), 0, 0, In_d, sin_d, cos_d, tan_d,
+    hipLaunchKernel(kernel_trig, dim3(LEN, 1, 1), dim3(1, 1, 1), 0, 0,
+                    In_d, sin_d, cos_d, tan_d,
                     sin_pd, cos_pd);
     hipMemcpy(sin_h, sin_d, SIZE, hipMemcpyDeviceToHost);
     hipMemcpy(cos_h, cos_d, SIZE, hipMemcpyDeviceToHost);
@@ -79,8 +86,28 @@ int main() {
     hipMemcpy(cos_ph, cos_pd, SIZE, hipMemcpyDeviceToHost);
     for (int i = 0; i < LEN; i++) {
         if (sin_h[i] != sin_ph[i] || cos_h[i] != cos_ph[i] || tan_h[i] * cos_h[i] != sin_h[i]) {
-            std::cout << "Failed!" << std::endl;
+            errors++;
+#if TEST_DEBUG
+            std::cout << "Check Failed!" << std::endl;
+            std::cout << " sin_h: " << sin_h[i] << " sin_ph: " << sin_ph[i] << "\n"
+                      << " cos_h: " << cos_h[i] << " cos_ph:" << cos_ph[i] << "\n"
+                      << " tan_h * cos_h: " << tan_h[i] * cos_h[i] << " sin_h[i]: " << sin_h[i] << "\n";
+#endif
         }
     }
-    passed();
+
+    HIP_ASSERT(hipFree(In_d));
+    HIP_ASSERT(hipFree(sin_d));
+    HIP_ASSERT(hipFree(cos_d));
+    HIP_ASSERT(hipFree(tan_d));
+    HIP_ASSERT(hipFree(sin_pd));
+    HIP_ASSERT(hipFree(cos_pd));
+
+    if (errors != 0) {
+        std::cout << "hip_trig FAILED!" << std::endl;
+        return -1;
+    } else {
+        std::cout << "hip_trig PASSED!" << std::endl;
+    }
+    return errors;
 }