intel / intel/llvm

Wrapping an Accessor in a Struct

Open
#9,434 2 comments 0 reactions 0 assignees View on GitHub
confirmed enhancement
Dominant language
LLVM
Stars
1.5k
Forks
854
Avg merge
3d 17h
Merged PRs (30d)
137

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:

```cpp
#include
#include

template
struct Accessor : AccT {
using Acc = AccT;

Accessor(Acc &&a) : AccT(std::move(a)) {}

/* non-functioning
template = 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(& (*this)[i][j]);
}
};

struct Mat : sycl::buffer {
Mat(size_t m, size_t n) : sycl::buffer({m, n}) {}
Mat(const sycl::buffer&F) : sycl::buffer(F) {}
};

// note: first arg. can be "struct mode_tag"
// second can be like "struct no_init : public detail::property"
template
auto host_accessor(Mat &X, Args ... args) {
return Accessor( sycl::host_accessor(X, args ...) );
}

template
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](https://github.com/ORNL/SYCLhash).

However, to avoid making the user write out all the types explicitly, the following (non-standards compliant hack) is required:
```
template class ctorType>
constexpr sycl::access_mode AccessMode_t(ctorType) {
return mode;
}

// template deduction helper
template
DeviceHash(Hash&,
sycl::handler&,
Descriptor d)
-> DeviceHash;
//-> DeviceHash; // hipSYCL

template
HostHash(Hash&, Descriptor d)
-> HostHash;
}
```

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.