From f22391c36286a3d02b075f710b6b44334886aed3 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri, 25 Oct 2019 11:14:24 +0100
Subject: [PATCH] Add missing operators, fix GCC compilation. (#1589)

[ROCm/clr commit: 40522e2b6a6714e0a6fb377b0b5b300843cb3e77]
---
 .../include/hip/hcc_detail/hip_vector_types.h | 42 ++++++++++++++++++-
 1 file changed, 41 insertions(+), 1 deletion(-)

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 f80745038a..b203d942a8 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
@@ -44,18 +44,35 @@ THE SOFTWARE.
         __attribute__((vector_size(__ROUND_UP_TO_NEXT_POT__(n) * sizeof(T))))
 #endif
 
-#if defined(__cplusplus)
+#if defined(__cplusplus) && defined(__clang__)
     #include <type_traits>
 
     namespace hip_impl {
         template<typename T, typename Vector, unsigned int idx>
         struct Scalar_accessor {
+            struct Address {
+                const Scalar_accessor* p;
+
+                __host__ __device__
+                operator const T*() const noexcept {
+                    return &reinterpret_cast<const T*>(p)[idx];
+                }
+                __host__ __device__
+                operator T*() noexcept {
+                    return &reinterpret_cast<T*>(
+                        const_cast<Scalar_accessor*>(p))[idx];
+                }
+            };
+
             // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor
             Vector data;
 
             __host__ __device__
             operator T() const noexcept { return data[idx]; }
 
+            __host__ __device__
+            Address operator&() const noexcept { return Address{this}; }
+
             __host__ __device__
             Scalar_accessor& operator=(T x) noexcept {
                 data[idx] = x;
@@ -63,6 +80,29 @@ THE SOFTWARE.
                 return *this;
             }
 
+            __host__ __device__
+            Scalar_accessor& operator++() noexcept {
+                ++data[idx];
+                return *this;
+            }
+            __host__ __device__
+            T operator++(int) noexcept {
+                auto r{data[idx]};
+                ++data[idx];
+                return *this;
+            }
+            __host__ __device__
+            Scalar_accessor& operator--() noexcept {
+                --data[idx];
+                return *this;
+            }
+            __host__ __device__
+            T operator--(int) noexcept {
+                auto r{data[idx]};
+                --data[idx];
+                return *this;
+            }
+
             __host__ __device__
             Scalar_accessor& operator+=(T x) noexcept {
                 data[idx] += x;