Merge 'master' into 'amd-master'

Change-Id: I3d27dfbda4db54c862dd2f243145ffdffc170179


[ROCm/clr commit: f1858aefb0]

projects/clr/hipamd/.gitignore

@@ -14,5 +14,5 @@ bin/hipBusBandwidth
 bin/hipDispatchLatency
 bin/hipify-clang
 include/hip/hcc_detail/hip_prof_str.h
-
+tags
 samples/1_Utils/hipInfo/hipInfo

projects/clr/hipamd/Jenkinsfile

@@ -168,7 +168,6 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
 
     // Cap the maximum amount of testing, in case of hangs
     // Excluding hipMultiThreadDevice-pyramid test from automation; due to its flakiness which requires some investigation
-    // Excluding hipLaunchParm test from automation, due to bug in HCC.
     timeout(time: 1, unit: 'HOURS')
     {
       stage("${platform} unit testing")
@@ -178,7 +177,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
             cd ${build_dir_rel}
             make install -j\$(nproc)
             make build_tests -i -j\$(nproc)
-            ctest -E "(pyramid|hipLaunchParm.tst)"
+            ctest -E pyramid
           """
         // If unit tests output a junit or xunit file in the future, jenkins can parse that file
         // to display test results on the dashboard

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h

@@ -206,12 +206,15 @@ class Coordinates {
 
     struct X {
         __device__ operator R() const { return f(0); }
+        __device__ uint32_t operator=(R _) { return f(0); }
     };
     struct Y {
         __device__ operator R() const { return f(1); }
+        __device__ uint32_t operator=(R _) { return f(1); }
     };
     struct Z {
         __device__ operator R() const { return f(2); }
+        __device__ uint32_t operator=(R _) { return f(2); }
     };
 
    public:

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -271,7 +271,7 @@ typedef struct dim3 {
     uint32_t y;  ///< y
     uint32_t z;  ///< z
 #ifdef __cplusplus
-    dim3(uint32_t _x = 1, uint32_t _y = 1, uint32_t _z = 1) : x(_x), y(_y), z(_z){};
+    __host__ __device__ dim3(uint32_t _x = 1, uint32_t _y = 1, uint32_t _z = 1) : x(_x), y(_y), z(_z){};
 #endif
 } dim3;
 

projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h

@@ -294,6 +294,7 @@ THE SOFTWARE.
             ++*this;
             return tmp;
         }
+
         inline __host__ __device__
         HIP_vector_type& operator--() noexcept
         {
@@ -306,12 +307,23 @@ THE SOFTWARE.
             --*this;
             return tmp;
         }
+
         inline __host__ __device__
         HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
         {
             data += x.data;
             return *this;
         }
+        template<
+            typename U,
+            typename std::enable_if<
+                std::is_convertible<U, T>{}>::type* = nullptr>
+        inline __host__ __device__
+        HIP_vector_type& operator+=(U x) noexcept
+        {
+            return *this += HIP_vector_type{x};
+        }
+
         inline __host__ __device__
         HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept
         {
@@ -327,18 +339,38 @@ THE SOFTWARE.
         {
             return *this -= HIP_vector_type{x};
         }
+
         inline __host__ __device__
         HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
         {
             data *= x.data;
             return *this;
         }
+        template<
+            typename U,
+            typename std::enable_if<
+                std::is_convertible<U, T>{}>::type* = nullptr>
+        inline __host__ __device__
+        HIP_vector_type& operator*=(U x) noexcept
+        {
+            return *this *= HIP_vector_type{x};
+        }
+
         inline __host__ __device__
         HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept
         {
             data /= x.data;
             return *this;
         }
+        template<
+            typename U,
+            typename std::enable_if<
+                std::is_convertible<U, T>{}>::type* = nullptr>
+        inline __host__ __device__
+        HIP_vector_type& operator/=(U x) noexcept
+        {
+            return *this /= HIP_vector_type{x};
+        }
 
         template<
             typename U = T,
@@ -361,6 +393,7 @@ THE SOFTWARE.
             r.data = ~r.data;
             return r;
         }
+
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
@@ -370,6 +403,7 @@ THE SOFTWARE.
             data %= x.data;
             return *this;
         }
+
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
@@ -379,6 +413,7 @@ THE SOFTWARE.
             data ^= x.data;
             return *this;
         }
+
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
@@ -388,6 +423,7 @@ THE SOFTWARE.
             data |= x.data;
             return *this;
         }
+
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
@@ -397,6 +433,7 @@ THE SOFTWARE.
             data &= x.data;
             return *this;
         }
+
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
@@ -406,6 +443,7 @@ THE SOFTWARE.
             data >>= x.data;
             return *this;
         }
+
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>