Compiler error when indexing sycl::accessor with dimensions > 1

[adaptivecpp]

Compiler error when indexing sycl::accessor with dimensions > 1

[illuhad]

Thanks for reporting. There are no unit tests yet to cover this slice API, so it's not yet really expected to work ;)
Personally, I would like to focus on fixing bugs and packaging in order to get the 0.8 release out as soon as possible. Since transitioning from "there's some code related to that feature" to "that feature is officially supported and routinely tested" is a bit more involved than fixing a bug, I would like to postpone that after the 0.8 release.

In the meantime, if feasible for your code base, I would propose to index with sycl::id<2>, as you did in the kernel, e.g.

png_image[sycl::id<2>{x,y}] = ...

[tomdeakin]

I'm also running into this issue.

It's very convenient for stencil codes to index the access as buf[x][y] += buf[x-1][y] + buf[x+1][y] + ....

It's tricky because the return type of this accessor is unspecified in the spec, which is probably sensible as it's an internal implementation detail. Ultimately when everything is specified you'd want this to come out as Equation 4.3, but I'm not totally convinced this is possible.

[illuhad]

So far I have no concerns that it is possible to implement equation 4.3 in that way (I assume this is just the linear id equation, I don't know the equation numbering from the top of my head ;) ). After all, the linear id equation is just the regular linearization equation that C and C++ use for constructs like array[x][y]. But I am curious to know why you are skeptical?

Is it an acceptable workaround for your stencil code as well to just use sycl::id for indexing?

[tomdeakin]

Sorry, yes that's the linear id equation!

Do we not have to worry about the intermediate accessor[x]. Can the compiler go directly from accessor[x][y] to accessor[x*n+y] by definition rather doing (accessor[x])[y] and hoping for the inliner to fix it up? The accessor isn't just a pointer type so it's really a operator[] that we have to define.

[illuhad]

Ah, so you are effectively asking about various implementation choices (compiler vs library).

I suppose this could probably also be implemented with some compiler trickery to replace the nested operator[] calls with a direct calculation of the linear id. But in hipSYCL, this was never really an option because it requires much, much more engineering on the compiler front (which is also much harder to maintain) and the performance benefit is unclear. As I am the only regular hipSYCL developer, I don't really have the resources to implement something in such an expensive way if it's not unavoidable for performance (I also don't know if the other implementations have decided that this is worth it? triSYCL seems to just rely on boost::multi_array_ref for this feature, so no compiler support either)

Because of this, the current implementation in hipSYCL (which is not yet working) returns a proxy object from operator[], which is actually just an accessor of dimension-1 that has a modified access offset to account for the offset of the index from the current dimension.
There is some overhead associated with this because the accessor class is very complex, so creating multiple accessors is not ideal and also a bit convoluted (which is probably why it's broken now) - back when this code was written the implementation was much simpler ;)
However, I don't think that an implementation with a proxy object necessarily has to be inferior. For example, the proxy object could just store the pointer to the actual accessor and collect the access indices in an id<dim>. Once the id is complete, it could just invoke the regular accessor operator[id<dim>]. I don't think this will be any slower than an implementation relying on compiler magic.
Thinking about it, I suspect that this would also be a much cleaner solution compared to the current implementation attempt, so rather than fixing the current code maybe we should just switch to this approach...

[illuhad]

Houston, we have a fix! I took some time yesterday evening to implement the idea outlined above in PR #186.

[illuhad]

Fixed by PR #186.

[tomdeakin]

Thanks!!