From ee5097f2c2c257c8dcd2aa67a003b84fb0e50dd2 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri, 1 Nov 2019 22:18:01 +0200
Subject: [PATCH] Accessors should work even when oddly volatile.

---
 include/hip/hcc_detail/hip_vector_types.h | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h
index b203d942a8..9764bc2a16 100644
--- a/include/hip/hcc_detail/hip_vector_types.h
+++ b/include/hip/hcc_detail/hip_vector_types.h
@@ -54,11 +54,11 @@ THE SOFTWARE.
                 const Scalar_accessor* p;
 
                 __host__ __device__
-                operator const T*() const noexcept {
+                operator const T*() const volatile noexcept {
                     return &reinterpret_cast<const T*>(p)[idx];
                 }
                 __host__ __device__
-                operator T*() noexcept {
+                operator T*() volatile noexcept {
                     return &reinterpret_cast<T*>(
                         const_cast<Scalar_accessor*>(p))[idx];
                 }
@@ -68,7 +68,7 @@ THE SOFTWARE.
             Vector data;
 
             __host__ __device__
-            operator T() const noexcept { return data[idx]; }
+            operator T() const volatile noexcept { return data[idx]; }
 
             __host__ __device__
             Address operator&() const noexcept { return Address{this}; }