Merge 'master' into 'amd-master'

Change-Id: I4784a1bb118a6000aeff23a90befec6abbceacae

bin/hipcc

@@ -114,14 +114,19 @@ if ($HIP_PLATFORM eq "clang") {
     }
     $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
     $HIPCC="$HIP_CLANG_PATH/clang++";
-    $HIPCXXFLAGS .= "-std=c++11 -I$HIP_PATH/include";
-    $HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc";
 
     # If $HIPCC clang++ is not compiled, use clang instead
     if ( ! -e $HIPCC ) {
         $HIPCC="$HIP_CLANG_PATH/clang";
-        $HIPLDFLAGS = "--driver-mode=g++ " . $HIPLDFLAGS;
+        $HIPLDFLAGS = "--driver-mode=g++";
     }
+
+    $HIP_CLANG_VERSION = `$HIPCC --version`;
+    $HIP_CLANG_VERSION=~/.*clang version ([^ ]+).*/;
+    $HIP_CLANG_VERSION=$1;
+
+    $HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/include -I$HIP_PATH/include";
+    $HIPLDFLAGS .= " --hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc";
 } elsif ($HIP_PLATFORM eq "hcc") {
     $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
 
@@ -439,7 +444,7 @@ foreach $arg (@ARGV)
                 $toolArgs .= " -x hip"
             }
         }
-        if (($arg =~ /\.cu$/) or ($arg =~ /\.cuh$/)) {
+        if (($arg =~ /\.cu$/) or ($arg =~ /\.cuh$/) or ($arg =~ /\.hip$/)) {
             $hasCU = 1;
             $needCXXFLAGS = 1;
             if ($HIP_PLATFORM eq 'clang') {

include/hip/hcc_detail/device_functions.h

@@ -30,6 +30,7 @@ THE SOFTWARE.
 #include <hip/hip_vector_types.h>
 #include <hip/hcc_detail/device_library_decls.h>
 #include <hip/hcc_detail/llvm_intrinsics.h>
+#include <stddef.h>
 
 typedef unsigned long ulong;
 typedef unsigned int uint;
@@ -970,4 +971,62 @@ unsigned __smid(void)
 #endif //defined(__clang__) && defined(__HIP__)
 
 
+// loop unrolling
+static inline __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size) {
+    auto dstPtr = static_cast<unsigned char*>(dst);
+    auto srcPtr = static_cast<const unsigned char*>(src);
+
+    while (size >= 4u) {
+        dstPtr[0] = srcPtr[0];
+        dstPtr[1] = srcPtr[1];
+        dstPtr[2] = srcPtr[2];
+        dstPtr[3] = srcPtr[3];
+
+        size -= 4u;
+        srcPtr += 4u;
+        dstPtr += 4u;
+    }
+    switch (size) {
+        case 3:
+            dstPtr[2] = srcPtr[2];
+        case 2:
+            dstPtr[1] = srcPtr[1];
+        case 1:
+            dstPtr[0] = srcPtr[0];
+    }
+
+    return dst;
+}
+
+static inline __device__ void* __hip_hc_memset(void* dst, unsigned char val, size_t size) {
+    auto dstPtr = static_cast<unsigned char*>(dst);
+
+    while (size >= 4u) {
+        dstPtr[0] = val;
+        dstPtr[1] = val;
+        dstPtr[2] = val;
+        dstPtr[3] = val;
+
+        size -= 4u;
+        dstPtr += 4u;
+    }
+    switch (size) {
+        case 3:
+            dstPtr[2] = val;
+        case 2:
+            dstPtr[1] = val;
+        case 1:
+            dstPtr[0] = val;
+    }
+
+    return dst;
+}
+static inline __device__ void* memcpy(void* dst, const void* src, size_t size) {
+    return __hip_hc_memcpy(dst, src, size);
+}
+
+static inline __device__ void* memset(void* ptr, int val, size_t size) {
+    unsigned char val8 = static_cast<unsigned char>(val);
+    return __hip_hc_memset(ptr, val8, size);
+}
 #endif

include/hip/hcc_detail/hip_runtime.h

@@ -260,25 +260,12 @@ static constexpr Coordinates<hc_get_workitem_id> threadIdx;
 
 #endif // defined __HCC__
 #if __HCC_OR_HIP_CLANG__
-extern "C" __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size);
-extern "C" __device__ void* __hip_hc_memset(void* ptr, uint8_t val, size_t size);
 extern "C" __device__ void* __hip_hc_malloc(size_t);
 extern "C" __device__ void* __hip_hc_free(void* ptr);
 
 static inline __device__ void* malloc(size_t size) { return __hip_hc_malloc(size); }
-
 static inline __device__ void* free(void* ptr) { return __hip_hc_free(ptr); }
 
-static inline __device__ void* memcpy(void* dst, const void* src, size_t size) {
-    return __hip_hc_memcpy(dst, src, size);
-}
-
-static inline __device__ void* memset(void* ptr, int val, size_t size) {
-    uint8_t val8 = static_cast<uint8_t>(val);
-    return __hip_hc_memset(ptr, val8, size);
-}
-
-
 #ifdef __HCC_ACCELERATOR__
 
 #ifdef HC_FEATURE_PRINTF
@@ -453,6 +440,16 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
 
 #include <hip/hcc_detail/math_functions.h>
 
+// Support std::complex.
+#pragma push_macro("__CUDA__")
+#define __CUDA__
+#include <__clang_cuda_math_forward_declares.h>
+#include <__clang_cuda_complex_builtins.h>
+#include <cuda_wrappers/algorithm>
+#include <cuda_wrappers/complex>
+#undef __CUDA__
+#pragma pop_macro("__CUDA__")
+
 #endif
 
 #endif  // HIP_HCC_DETAIL_RUNTIME_H

include/hip/hcc_detail/hip_vector_types.h

@@ -120,7 +120,8 @@ THE SOFTWARE.
         }
         template< // TODO: constrain based on type as well.
             typename... Us,
-            typename std::enable_if<sizeof...(Us) == rank>::type* = nullptr>
+            typename std::enable_if<
+                (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
         __host__ __device__
         HIP_vector_type(Us... xs) noexcept { data = Native_vec_{xs...}; }
         __host__ __device__
@@ -769,4 +770,4 @@ DECLOP_MAKE_TWO_COMPONENT(signed long, longlong2);
 DECLOP_MAKE_THREE_COMPONENT(signed long, longlong3);
 DECLOP_MAKE_FOUR_COMPONENT(signed long, longlong4);
 
-#endif
+#endif

src/device_util.cpp

@@ -92,58 +92,6 @@ __device__ void* __hip_hc_free(void* ptr) {
     return nullptr;
 }
 
-
-// loop unrolling
-__device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size) {
-    auto dstPtr = static_cast<uint8_t*>(dst);
-    auto srcPtr = static_cast<const uint8_t*>(src);
-
-    while (size >= 4u) {
-        dstPtr[0] = srcPtr[0];
-        dstPtr[1] = srcPtr[1];
-        dstPtr[2] = srcPtr[2];
-        dstPtr[3] = srcPtr[3];
-
-        size -= 4u;
-        srcPtr += 4u;
-        dstPtr += 4u;
-    }
-    switch (size) {
-        case 3:
-            dstPtr[2] = srcPtr[2];
-        case 2:
-            dstPtr[1] = srcPtr[1];
-        case 1:
-            dstPtr[0] = srcPtr[0];
-    }
-
-    return dst;
-}
-
-__device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) {
-    auto dstPtr = static_cast<uint8_t*>(dst);
-
-    while (size >= 4u) {
-        dstPtr[0] = val;
-        dstPtr[1] = val;
-        dstPtr[2] = val;
-        dstPtr[3] = val;
-
-        size -= 4u;
-        dstPtr += 4u;
-    }
-    switch (size) {
-        case 3:
-            dstPtr[2] = val;
-        case 2:
-            dstPtr[1] = val;
-        case 1:
-            dstPtr[0] = val;
-    }
-
-    return dst;
-}
-
 // abort
 __device__ void abort() { return hc::abort(); }
 

src/functional_grid_launch.inl

@@ -123,21 +123,17 @@ namespace hip_impl
             };
         }
 
-        for (auto&& agent_kernel : it0->second) {
-            if (agent.handle == agent_kernel.first.handle) {
-                hipModuleLaunchKernel(
-                    agent_kernel.second,
-                    numBlocks.x,
-                    numBlocks.y,
-                    numBlocks.z,
-                    dimBlocks.x,
-                    dimBlocks.y,
-                    dimBlocks.z,
-                    sharedMemBytes,
-                    stream,
-                    nullptr,
-                    kernarg);
-            }
-        }
+        hipModuleLaunchKernel(
+            it1->second,
+            numBlocks.x,
+            numBlocks.y,
+            numBlocks.z,
+            dimBlocks.x,
+            dimBlocks.y,
+            dimBlocks.z,
+            sharedMemBytes,
+            stream,
+            nullptr,
+            kernarg);
     }
 }

tests/hipify-clang/dynamic_shared_memory.cu

@@ -0,0 +1,44 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// Taken from Jonathan Hui blog https://jhui.github.io/2017/03/06/CUDA
+
+#include <stdio.h>
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda.h>
+
+__global__ void dynamicReverse(int *d, int n)
+{
+  // Dynamic shared memory
+  // CHECK: HIP_DYNAMIC_SHARED(int, s);
+  extern __shared__ int s[];
+  int t = threadIdx.x;
+  int tr = n-t-1;
+  s[t] = d[t];
+  __syncthreads();
+  d[t] = s[tr];
+}
+
+int main(void)
+{
+  const int n = 64;
+  int a[n], r[n], d[n];
+
+  for (int i = 0; i < n; i++) {
+    a[i] = i;
+    r[i] = n-i-1;
+    d[i] = 0;
+  }
+
+  int *d_d;
+  // CHECK: hipMalloc(&d_d, n * sizeof(int));
+  cudaMalloc(&d_d, n * sizeof(int));
+  // run version with dynamic shared memory
+  // CHECK: hipMemcpy(d_d, a, n*sizeof(int), hipMemcpyHostToDevice);
+  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
+  // CHECK: hipLaunchKernelGGL(dynamicReverse, dim3(1), dim3(n), n*sizeof(int), 0, d_d, n);
+  dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
+  // CHECK: hipMemcpy(d, d_d, n*sizeof(int), hipMemcpyDeviceToHost);
+  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
+  for (int i = 0; i < n; i++)
+    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)n", i, i, d[i], r[i]);
+}

tests/hipify-clang/static_shared_memory.cu

@@ -0,0 +1,43 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// Taken from Jonathan Hui blog https://jhui.github.io/2017/03/06/CUDA
+
+#include <stdio.h>
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda.h>
+
+__global__ void staticReverse(int *d, int n)
+{
+  __shared__ int s[64];
+  int t = threadIdx.x;
+  int tr = n-t-1;
+  s[t] = d[t];
+  // Will not conttinue until all threads completed.
+  __syncthreads();
+  d[t] = s[tr];
+}
+
+int main(void)
+{
+  const int n = 64;
+  int a[n], r[n], d[n];
+
+  for (int i = 0; i < n; i++) {
+    a[i] = i;
+    r[i] = n-i-1;
+    d[i] = 0;
+  }
+
+  int *d_d;
+  // CHECK: hipMalloc(&d_d, n * sizeof(int));
+  cudaMalloc(&d_d, n * sizeof(int));
+  // run version with static shared memory
+  // CHECK: hipMemcpy(d_d, a, n*sizeof(int), hipMemcpyHostToDevice);
+  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
+  // CHECK: hipLaunchKernelGGL(staticReverse, dim3(1), dim3(n), 0, 0, d_d, n);
+  staticReverse<<<1,n>>>(d_d, n);
+  // CHECK: hipMemcpy(d, d_d, n*sizeof(int), hipMemcpyDeviceToHost);
+  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
+  for (int i = 0; i < n; i++)
+    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)n", i, i, d[i], r[i]);
+}

tests/src/deviceLib/hipStdComplex.cpp

@@ -0,0 +1,161 @@
+/*
+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
+ * RUN: %t
+ * HIT_END
+ */
+#include "test_common.h"
+#include <iostream>
+#include <complex>
+
+// Tolerance for error
+const double tolerance = 1e-6;
+const bool verbose = false;
+
+#define LEN 64
+
+#define ALL_FUN \
+  OP(add) \
+  OP(sub) \
+  OP(mul) \
+  OP(div) \
+  OP(abs) \
+  OP(arg) \
+  OP(sin) \
+  OP(cos)
+
+#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.
+// If the function returns real number, converts it to a complex number.
+#define ONE_ARG(func) \
+  case CK_##func: \
+    return std::complex<FloatT>(std::func(A));
+
+template<typename FloatT>
+__device__ __host__ std::complex<FloatT> calc(std::complex<FloatT> A,
+                                              std::complex<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;
+    
+    ONE_ARG(abs)
+    ONE_ARG(arg)
+    ONE_ARG(sin)
+    ONE_ARG(cos)
+  }
+}
+
+template<typename FloatT>
+__global__ void kernel(hipLaunchParm lp, std::complex<FloatT>* A,
+                       std::complex<FloatT>* B, std::complex<FloatT>* C,
+                       enum CalcKind CK) {
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
+    C[tx] = calc<FloatT>(A[tx], B[tx], CK);
+}
+
+template<typename FloatT>
+void test() {
+    typedef std::complex<FloatT> ComplexT;
+    ComplexT *A, *Ad, *B, *Bd, *C, *Cd, *D;
+    A = new ComplexT[LEN];
+    B = new ComplexT[LEN];
+    C = new ComplexT[LEN];
+    D = new ComplexT[LEN];
+    hipMalloc((void**)&Ad, sizeof(ComplexT)*LEN);
+    hipMalloc((void**)&Bd, sizeof(ComplexT)*LEN);
+    hipMalloc((void**)&Cd, sizeof(ComplexT)*LEN);
+    
+    for (uint32_t i = 0; i < LEN; i++) {
+        A[i] = ComplexT((i + 1) * 1.0f, (i + 2) * 1.0f);
+        B[i] = A[i];
+        C[i] = A[i];
+    }
+    hipMemcpy(Ad, A, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice);
+    hipMemcpy(Bd, B, sizeof(ComplexT)*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) {
+      hipLaunchKernel(kernel<FloatT>, dim3(1), dim3(LEN), 0, 0, Ad, Bd, Cd, CK);
+      hipMemcpy(C, Cd, sizeof(ComplexT)*LEN, hipMemcpyDeviceToHost);
+      for (int i = 0; i < LEN; i++) {
+        ComplexT 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() {
+// ToDo: Fix bug in HCC causing linking error at -O0.
+#ifndef __HCC__
+  test<float>();
+  test<double>();
+#endif
+  passed();
+  return 0;
+}

tests/src/deviceLib/hipVectorTypes.cpp

@@ -38,18 +38,23 @@ THE SOFTWARE.
 
 using namespace std;
 
-bool integer_unary_tests(...) {
+template<
+    typename V,
+    Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+bool integer_unary_tests(V&, V&) {
     return true;
 }
 
-bool integer_binary_tests(...) {
+template<
+    typename V,
+    Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+bool integer_binary_tests(V&, V&, V&...) {
     return true;
 }
 
 template<
     typename V,
     Enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
-__device__
 bool integer_unary_tests(V& f1, V& f2) {
     f1 %= f2;
     if (f1 != V{0}) return false;
@@ -71,7 +76,6 @@ bool integer_unary_tests(V& f1, V& f2) {
 template<
     typename V,
     Enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
-__device__
 bool integer_binary_tests(V& f1, V& f2, V& f3) {
     f3 = f1 % f2;
     if (f3 != V{0}) return false;

tests/src/deviceLib/hipVectorTypesDevice.cpp

@@ -40,6 +40,7 @@ using namespace std;
 template<
     typename V,
     Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+__device__
 constexpr
 bool integer_unary_tests(const V&, const V&) {
     return true;
@@ -73,6 +74,7 @@ bool integer_unary_tests(V& f1, V& f2) {
 template<
     typename V,
     Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+__device__
 constexpr
 bool integer_binary_tests(const V&, const V&, const V&) {
     return true;