Description
This issue is not about a bug, but imperfect handling of bad user code.
Per both SYCL-1.2.1 and SYCL-2020 (provisional), the source accessor to handler::copy must be either read or read_write.
However, the code below compiles and runs fine with the mode set to discard_write or write.
#include <CL/sycl.hpp>
#include <vector>
static constexpr int size = 32;
int main() {
cl::sycl::device dev{cl::sycl::gpu_selector()};
cl::sycl::queue q{dev};
cl::sycl::buffer<int> d_x(cl::sycl::range<1>{size});
std::vector<int> h_x(size, 0);
q.submit([&](cl::sycl::handler &cgh) {
auto accessor = d_x.get_access<cl::sycl::access::mode::write>(cgh);
cgh.fill(accessor, 1);
});
q.submit([&](cl::sycl::handler &cgh) {
auto accessor = d_x.get_access<cl::sycl::access::mode::write>(cgh); // <<<
cgh.copy(accessor, h_x.data());
});
q.wait_and_throw();
return 0;
}An example where such a diagnostic would have been helpful (besides just being pedantic):
A kernel that only reads the buffer is enqueued after the cgh.copy. If the code had the correct access mode for copying, it would have been possible to execute the kernel and copy the data simultaneously. With the wrong access mode as in the example above, the code works, but the kernel is scheduled to wait for the data copy to complete. A simple static_assert here would have warned the user of their mistake and of missed optimization opportunity.