ROCm
Missing volatile assignment operator in half type
Hi,
In rocPRIM, I would like to add a specialised template function in rocPRIM that deals with volatile loading/storing, e.g.
__device__ inline
half load_volatile(half * input)
{
half retval;
auto input_ptr = reinterpret_cast<volatile half *>(input);
retval = *input_ptr;
return retval;
}
The reason for this is to improve the performance of casting; a reinterpret_cast<volatile short *> leads to poor performance, while reinterpret_cast<volatile _Float16 *> seems to have issues compiling with Caffe2.
I tried to simply cast reinterpret_cast<volatile half *> as you would with a fundamental type (though with a fundamental type, there's no need for reinterpret_cast, and you can use const_cast), but I'm met with the following build error
half_case.cpp:65:14: error: no viable conversion from 'volatile __half' to 'const __half_raw'
retval = *input_ptr;
^~~~~~~~~~
./half_case.h:12:8: note: candidate constructor (the implicit copy constructor) not viable: no known conversion from 'volatile __half' to 'const __half_raw &' for 1st argument
struct __half_raw {
^
./half_case.h:12:8: note: candidate constructor (the implicit move constructor) not viable: no known conversion from 'volatile __half' to '__half_raw &&' for 1st argument
./half_case.h:197:5: note: candidate function not viable: 'this' argument has type 'volatile __half', but method is not marked volatile
operator __half_raw() const { return __half_raw{data}; }
^
./half_case.h:27:45: note: candidate template ignored: disabled by 'enable_if' [with T = __half_raw]
using Enable_if_t = typename std::enable_if<cond, T>::type;
^
./half_case.h:27:45: note: candidate template ignored: disabled by 'enable_if' [with T = __half_raw]
./half_case.h:199:5: note: candidate function
operator volatile __half_raw() const volatile
^
./half_case.h:74:41: note: passing argument to parameter 'x' here
__half& operator=(const __half_raw& x)
Why is there an assignment operator for volatile __half_raw but not for volatile __half?
@AJcodes Apologies for the lack of response. Do you still need assistance with this ticket? Thanks!
