Separate volatile for clarity. Handle assignment.

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<const T*>(p)[idx];
+                }
                 __host__ __device__
                 operator const T*() const volatile noexcept {
                     return &reinterpret_cast<const T*>(p)[idx];
                 }
                 __host__ __device__
+                operator T*() noexcept {
+                    return &reinterpret_cast<T*>(
+                        const_cast<Scalar_accessor*>(p))[idx];
+                }
+                __host__ __device__
                 operator T*() volatile noexcept {
                     return &reinterpret_cast<T*>(
                         const_cast<Scalar_accessor*>(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 {