Merge 'master' into 'amd-master'

Change-Id: Id88d6b69117acd7c73a1279bbb4a5b1a5815b2e5 [ROCm/clr commit: eec176e474]
2018-09-18 04:09:49 -05:00
parent 73c9b996fb 814ac65527
commit d1f4a50ee9
10 changed files with 23 additions and 33 deletions
@@ -29,7 +29,7 @@ THE SOFTWARE.
    #include <utility>
 #endif

-#if defined(__clang__) && (__clang_major__ > 3)
+#if defined(__clang__) && (__clang_major__ > 5)
    typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2)));

    struct __half_raw {
@@ -39,6 +39,8 @@ THE SOFTWARE.

 #define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE)

+#if __HCC__ || __HIP__
+
 __attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
 __attribute__((weak)) __device__
    uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
@@ -102,4 +104,6 @@ extern "C" inline __device__ void* __hip_free(void* ptr) {
    return nullptr;
 }

+#endif
+
 #endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H
@@ -334,13 +334,13 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
 typedef int hipLaunchParm;

 template <typename... Args, typename F = void (*)(Args...)>
-inline void hipLaunchKernelGGL(F kernelName, const dim3& numblocks, const dim3& numthreads,
+inline void hipLaunchKernelGGL(F&& kernelName, const dim3& numblocks, const dim3& numthreads,
                               unsigned memperblock, hipStream_t streamId, Args... args) {
  kernelName<<<numblocks, numthreads, memperblock, streamId>>>(args...);
 }

 template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
-inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
+inline void hipLaunchKernel(F&& kernel, const dim3& numBlocks, const dim3& dimBlocks,
                            std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
    hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
                       std::move(args)...);
@@ -118,6 +118,9 @@ typedef int hipLaunchParm;
    }
 #endif

+#define __clock() clock()
+#define __clock64() clock64()
+
 #endif

 #endif
@@ -25,7 +25,7 @@ popd () {
 function setupENV()
 {
    sudo apt-get update
-    sudo apt-get install dpkg-dev rpm doxygen libelf-dev
+    sudo apt-get install dpkg-dev rpm doxygen libelf-dev rename
 }

 function buildHIP()
@@ -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);
-}
+}
@@ -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();
 }
@@ -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"

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