From 1a3d02fb08df508c0b4fced138ad8bb089ce64c9 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/clr commit: ed0d6ec51e858e380cf2bd07ff8add0c0940348a] --- .../include/hip/hcc_detail/hip_vector_types.h | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index 9764bc2a16..75c79d422c 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/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 {