From e87f9b1944506368d743734ba2f989510298ca39 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Nov 2019 22:02:08 +0200 Subject: [PATCH] Separate volatile for clarity. Handle assignment. [ROCm/hip commit: 99b9d5449fcfc5e4ab1d23c2992dd7f4187e9127] --- .../include/hip/hcc_detail/hip_vector_types.h | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 9764bc2a16..75c79d422c 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -53,11 +53,20 @@ THE SOFTWARE. struct Address { const Scalar_accessor* p; + __host__ __device__ + operator const T*() const noexcept { + return &reinterpret_cast(p)[idx]; + } __host__ __device__ operator const T*() const volatile noexcept { return &reinterpret_cast(p)[idx]; } __host__ __device__ + operator T*() noexcept { + return &reinterpret_cast( + const_cast(p))[idx]; + } + __host__ __device__ operator T*() volatile noexcept { return &reinterpret_cast( const_cast(p))[idx]; @@ -67,6 +76,8 @@ THE SOFTWARE. // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor Vector data; + __host__ __device__ + operator T() const noexcept { return data[idx]; } __host__ __device__ operator T() const volatile noexcept { return data[idx]; } @@ -79,6 +90,12 @@ THE SOFTWARE. return *this; } + __host__ __device__ + volatile Scalar_accessor& operator=(T x) volatile noexcept { + data[idx] = x; + + return *this; + } __host__ __device__ Scalar_accessor& operator++() noexcept {