[SYCL][CUDA] assert when buffer size is 0 ?

[zjin-lcf]

When the size of a device buffer is zero, the runtime will show assertion error. Could the assertion occur only when the size is negative ?

  what():  SYCL buffer size is zero. To create a device accessor, SYCL buffer size must be greater than zero. -30 

In CUDA, cudaMalloc() does not return an error when the size is zero.

Porting a CUDA program to SYCL using SYCL buffers is smooth when a buffer size of zero does not cause the assert error.

[AerialMantis]

So zero-sized buffers result in an error as it was previously the case in SYCL 2020 that zero-sized buffers were not permitted, even if the backend could allow it.

However, the SYCL 2020 spec was recently amended to permit zero-sized buffers:

Zero-sized buffers are permitted. In this case, the buffer still handles ownership as normal, but does not
need to store the zero-sized data itself.

As this is now permitted by the spec I think it's a valid enhancement, and I believe this is a general issue which would apply to all backends, so not just for the CUDA backend.

cc @bader

[AerialMantis]

I believe that is the intention of the recent clarification to the specification, the idea being that creating an accessor to a zero-sized buffer would still represent a data dependency but no copy would be performed as there is no data to copy. I think we should change this so that creating an accessor from a zero-sized buffer is valid to conform with the change to the specification.

[alexbatashev]

What's the use case for buffers and accessors anyway? You can't really use them, so, this sounds more like a great opportunity to hurt yourself in a hard-to-catch manner. Right now buffer clearly indicates to users what's wrong, eliminating need for bounds checking all over application.

[zjin-lcf]

Please take a look at https://github.com/zjin-lcf/HeCBench/tree/master/mnist-sycl
I am not sure if you think this is a good example to show the use case.
As a workaround, copy is disabled when the copy size is 0 in my example.

[AerialMantis]

So I believe the use case for this was mainly for writing generic code, where the size may or may not be zero, @alexbatashev you're right this does introduce a requirement on the user to at some point check to avoid dereferencing the accessor as this would be undefined behaviour.

[zjin-lcf]

Running this example would also also segfault without the workaround proposed in #4960
My comment is not relevant to the title of the topic, but I feel people need to be aware of it. Thanks.