Missing const T& conversion for Scalar_accessor

[Oblomov]

Hello,

I'm in the process of adding HIP support to GPUSPH, and one of the obstacles I've encountered so far is the lack of a operator const T&() const conversion operator for Scalar_accessor. The issue manifests only when building the host code using g++ $(hipconfig -C), and can be observed e.g. with this simple test case:

#include <hip/hip_runtime.h>

static __device__ __forceinline__ __host__ float clamp(const float &f, const float &a, const float &b)
{
	return fmaxf(a, fminf(f, b));
}

static __device__ __forceinline__ __host__ float2 clamp(const float2 &f, const float &a, const float &b)
{
	return make_float2( clamp(f.x, a, b), clamp(f.y, a, b) );
}

compiled with g++ $(hipconfig -C | tr -d \") -c test.cc, giving the error

test.cc: In function ‘float2 clamp(const float2&, const float&, const float&)’:
test.cc:10:30: error: invalid user-defined conversion from ‘const hip_impl::Scalar_accessor<float, float [2], 0>’ to ‘const float&’ [-fpermissive]
   10 |  return make_float2( clamp(f.x, a, b), clamp(f.y, a, b) );
      |                            ~~^
In file included from /opt/rocm-4.3.0/hip/include/hip/amd_detail/channel_descriptor.h:28,
                 from /opt/rocm-4.3.0/hip/include/hip/amd_detail/hip_texture_types.h:38,
                 from /opt/rocm-4.3.0/hip/include/hip/amd_detail/hip_runtime_api.h:40,
                 from /opt/rocm-4.3.0/hip/include/hip/hip_runtime_api.h:410,
                 from /opt/rocm-4.3.0/hip/include/hip/hip_runtime.h:113,
                 from test.cc:1:
/opt/rocm-4.3.0/hip/include/hip/amd_detail/hip_vector_types.h:141:13: note: candidate is: ‘hip_impl::Scalar_accessor< <template-parameter-1-1>, <template-parameter-1-2>, <anonymous> >::operator T&() [with T = float; Vector = float [2]; unsigned int idx = 0]’ (near match)
  141 |             operator T&() noexcept {
      |             ^~~~~~~~

The same error does not appear when compiling with hipcc, or when using clang++ instead of g++ as host compiler, so it's possible that the issue is due to the different representation used for vector types in the two cases. I noticed that the Scalar_accessor does provide operator T() const conversions, but these aren't apparently sufficient to satisfy g++. A solution could be to conver them into const T& operators, that should work just as well for both the T and const T& case.

[Oblomov]

Another example, with a simpler test-case that also fails with clang++:

#include <hip/hip_runtime.h>

typedef ushort4 particleinfo;

static __forceinline__ __host__ __device__ __attribute__((pure))
const ushort& type(const particleinfo &info)
{ return info.x; }

compiled with clang++ $(hipconfig -C) -c will give the warning:

test2.cc:7:10: warning: returning reference to local temporary object [-Wreturn-stack-address]
{ return info.x; }

Replacing the operator T() const with appropriate operator const T&() const operators fixes this issue as well.

[ppanchad-amd]

@Oblomov Apologies for the lack of response. Can you please test with latest ROCm 6.0.2 (HIP 6.0.32831)? If resolved, please close ticket. Thanks!

[Oblomov]

Thank you, this seems to be fixed in 6.0.2