diff --git a/projects/rocrand/library/include/rocrand/rocrand_common.h b/projects/rocrand/library/include/rocrand/rocrand_common.h index 7ca96ec48212..e4c3df65ebc4 100644 --- a/projects/rocrand/library/include/rocrand/rocrand_common.h +++ b/projects/rocrand/library/include/rocrand/rocrand_common.h @@ -87,6 +87,32 @@ #define ROCRAND_DEPRECATED(msg) #endif +// This is an accessor macro for HIP vector types (eg. int2, float4, etc.). +// Prior to HIP 7.0, individual elements could be accessed through the +// data member: +// +// int2 vec; +// vec.data[0] = 1; +// +// Beginning with HIP 7.0, the data member is hidden, and individual +// elements must be accessed like this: +// +// int2 vec; +// vec[0] = 1; +// +// You can use the macro like this: +// +// int2 vec; +// ROCRAND_HIPVEC_ACCESS(vec)[0]; +// +#if defined(__HIP_PLATFORM_AMD__) + #if HIP_VERSION_MAJOR < 7 + #define ROCRAND_HIPVEC_ACCESS(x) x.data + #else + #define ROCRAND_HIPVEC_ACCESS(x) x + #endif +#endif + namespace rocrand_device { namespace detail { diff --git a/projects/rocrand/library/include/rocrand/rocrand_philox4x32_10.h b/projects/rocrand/library/include/rocrand/rocrand_philox4x32_10.h index 02b84a1f0678..022e8a2f389d 100644 --- a/projects/rocrand/library/include/rocrand/rocrand_philox4x32_10.h +++ b/projects/rocrand/library/include/rocrand/rocrand_philox4x32_10.h @@ -172,7 +172,7 @@ class philox4x32_10_engine __forceinline__ __device__ __host__ unsigned int next() { #if defined(__HIP_PLATFORM_AMD__) - unsigned int ret = m_state.result.data[m_state.substate]; + unsigned int ret = ROCRAND_HIPVEC_ACCESS(m_state.result)[m_state.substate]; #else unsigned int ret = (&m_state.result.x)[m_state.substate]; #endif diff --git a/projects/rocrand/library/include/rocrand/rocrand_threefry2_impl.h b/projects/rocrand/library/include/rocrand/rocrand_threefry2_impl.h index 54d985549bc1..97ee5ffc2bc9 100644 --- a/projects/rocrand/library/include/rocrand/rocrand_threefry2_impl.h +++ b/projects/rocrand/library/include/rocrand/rocrand_threefry2_impl.h @@ -138,7 +138,7 @@ class threefry_engine2_base __forceinline__ __device__ __host__ value next() { #if defined(__HIP_PLATFORM_AMD__) - value ret = m_state.result.data[m_state.substate]; + value ret = ROCRAND_HIPVEC_ACCESS(m_state.result)[m_state.substate]; #else value ret = (&m_state.result.x)[m_state.substate]; #endif diff --git a/projects/rocrand/library/include/rocrand/rocrand_threefry4_impl.h b/projects/rocrand/library/include/rocrand/rocrand_threefry4_impl.h index 99587058c8d7..4c93216e98c8 100644 --- a/projects/rocrand/library/include/rocrand/rocrand_threefry4_impl.h +++ b/projects/rocrand/library/include/rocrand/rocrand_threefry4_impl.h @@ -153,7 +153,7 @@ class threefry_engine4_base __forceinline__ __device__ __host__ value next() { #if defined(__HIP_PLATFORM_AMD__) - value ret = m_state.result.data[m_state.substate]; + value ret = ROCRAND_HIPVEC_ACCESS(m_state.result)[m_state.substate]; #else value ret = (&m_state.result.x)[m_state.substate]; #endif