Accessor equality operator broken in device code

SYCL development
1

Hi, looks like the accessor::operator== is broken in device code as of 1.3.0. This example

  cl::sycl::buffer<size_t, 1> buf(cl::sycl::range<1>(8));
  queue.submit([&](cl::sycl::handler& cgh) {
    auto acc = buf.get_access<cl::sycl::access::mode::discard_write>(cgh);
    cgh.parallel_for<class Foo>(buf.get_range(), [=](cl::sycl::id<1> id) {
        if(acc == acc) acc[id] = 1337;
    });
  });

causes the error

 error: [Computecpp:CC0006]: virtual function call is not allowed within SYCL device code
2

Woops, I just realized that other than most “common reference semantics” the equality operator only has to be defined on the host! Sorry for the noise :slight_smile:

3

Hi @psalz,

Did this work in previous versions? I’m trying to think what might have made it break because I don’t remember such a change going in (though my memory is very imperfect). What use case did you have for this? It seems slightly odd to me to think of how I might use this in a kernel, but I’m sure I’ve overlooked something. Would be great if you can share this!

I think that the operation should be possible to implement on device, so maybe there’s a specification suggestion in this issue.

Duncan.

4

Hi,

Did this work in previous versions?

I don’t think so, I just tried it with 1.0.0 and I’m getting the same error!

What use case did you have for this? It seems slightly odd to me to think of how I might use this in a kernel, but I’m sure I’ve overlooked something.

Haha yeah no I don’t have a real use case, I’m writing a custom wrapper around accessors and wanted to provide compatibility with as much of the original interface as possible.

I think that the operation should be possible to implement on device, so maybe there’s a specification suggestion in this issue.

Yes it seems weird to me that this is not already part of the spec, but maybe it really was just that nobody could come up with a compelling use case for it. As I’ve been opening too many issues on the spec repo lately I think I’ll hold off on this one for now ;-).

5

Thanks for your answers. I guess I originally read “as of 1.3.0” as meaning it worked before, which was a bit surprising since nothing I could see changed!

For keeping the same interface, maybe you could compare the pointers (through get_pointer) and then the sizes/ranges etc.? Depends what the original interface’s notion of equality is. Comparing two pointers has the benefit of being exceptionally cheap, even if it is missing the other bits (ranges, offsets, …). I’m sure the group doesn’t mind more suggestions, everything that can make the spec better is welcome I’m sure!