Wrapping an Accessor in a Struct #9434
Description
Describe the feature
Please change this issue tag to "feature / discussion".
I'd like to implement my own device-side containers that wrap accessors using code such as:
#include <sycl/sycl.hpp>
#include <stdio.h>
template <typename AccT>
struct Accessor : AccT {
using Acc = AccT;
Accessor(Acc &&a) : AccT(std::move(a)) {}
/* non-functioning
template <bool use=true, std::enable_if_t< use &&
Acc::Mode != sycl::access::mode::read, bool> = true >
float &operator()(int i, int j) {
return (*this)[i][j];
}*/
// Hack because it's not possible to determine whether
// AccT has read_only access mode?
float &operator()(int i, int j) const {
return *const_cast<float *>(& (*this)[i][j]);
}
};
struct Mat : sycl::buffer<float,2> {
Mat(size_t m, size_t n) : sycl::buffer<float,2>({m, n}) {}
Mat(const sycl::buffer<float,2>&F) : sycl::buffer<float,2>(F) {}
};
// note: first arg. can be "struct mode_tag<Descriptor>"
// second can be like "struct no_init : public detail::property"
template <typename ...Args>
auto host_accessor(Mat &X, Args ... args) {
return Accessor( sycl::host_accessor(X, args ...) );
}
template <typename ...Args>
auto accessor(sycl::handler &cgh, Mat &X, Args ... args) {
sycl::accessor A(X, cgh, args...);
return Accessor( std::move(A) );
}
int main(int argc, char *argv[]) {
float err = 0.0f;
sycl::queue queue;
Mat x(20,20);
Mat y(x);
queue.submit([&](sycl::handler &cgh) {
auto X = accessor(cgh, x, sycl::write_only, sycl::no_init);
cgh.parallel_for(sycl::range(20), [=](sycl::id<1> idx) {
for(int j=0; j<20; j++) {
X(idx,j) = j;
}
});
});
{
auto X = host_accessor(y, sycl::read_only);
for(int i=0; i<20; i++) {
printf("%f\n", X(i,i));
err += X(i,i) > i ? X(i,i)-i : i-X(i,i);
}
}
return err > 1e-5;
}But, the templated nature of SYCL accessors does not to provide an easy way to wrap an accessor inside a struct. Sophisticated user-defined accessor objects may require this functionality. The wrapping issue stems mainly from sycl::accessor not re-exporting its Mode tag as a public typedef. Alternately, the SYCL spec could be fixed here: inline constexpr __unspecified__ read_only;. hipSYCL's solution to defining mode tags is really great for this. Probably an accessor re-exporting its mode tag would be best though.
Additional context
Adding more template parameters to the wrapper is one solution -- which I have done here.
However, to avoid making the user write out all the types explicitly, the following (non-standards compliant hack) is required:
template <sycl::access_mode mode, template<sycl::access_mode> class ctorType>
constexpr sycl::access_mode AccessMode_t(ctorType<mode>) {
return mode;
}
// template deduction helper
template<typename T, int width, class Descriptor>
DeviceHash(Hash<T,width>&,
sycl::handler&,
Descriptor d)
-> DeviceHash<T, width, AccessMode_t(d), sycl::target::device >;
//-> DeviceHash<T, width, Descriptor::mode, Descriptor::target>; // hipSYCL
template<typename T, int width, class Descriptor>
HostHash(Hash<T,width>&, Descriptor d)
-> HostHash<T, width, AccessMode_t(d) >;
}