diff --git a/.gitignore b/.gitignore index 3c94b6aaa0..c4ed31eb6d 100644 --- a/.gitignore +++ b/.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 diff --git a/Jenkinsfile b/Jenkinsfile index e583ff5019..89870edeee 100644 --- a/Jenkinsfile +++ b/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 diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 6ac402a6f7..c4ef1d829f 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/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: diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index e70f4954db..97df1f4219 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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; diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 953e4eb5f8..3407b8a752 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/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{}>::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{}>::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{}>::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{}>::type* = nullptr> @@ -370,6 +403,7 @@ THE SOFTWARE. data %= x.data; return *this; } + template< typename U = T, typename std::enable_if{}>::type* = nullptr> @@ -379,6 +413,7 @@ THE SOFTWARE. data ^= x.data; return *this; } + template< typename U = T, typename std::enable_if{}>::type* = nullptr> @@ -388,6 +423,7 @@ THE SOFTWARE. data |= x.data; return *this; } + template< typename U = T, typename std::enable_if{}>::type* = nullptr> @@ -397,6 +433,7 @@ THE SOFTWARE. data &= x.data; return *this; } + template< typename U = T, typename std::enable_if{}>::type* = nullptr> @@ -406,6 +443,7 @@ THE SOFTWARE. data >>= x.data; return *this; } + template< typename U = T, typename std::enable_if{}>::type* = nullptr>