[SYCL] Ensure correct access mode in handler::copy (#3109)

[al42and]

This patch introduces compile-time checks to cl::sycl::handler::copy functions that ensure the correctness of accessors' access mode (section 4.8.6 of SYCL-1.2.1 specification).

Closes #3109

[al42and]

Unfortunately, I don't have access to Jenkins (Server Not Found). Could you please share the error message of failing Jenkins/Precommit check?

[AlexeySachkov]

Unfortunately, I don't have access to Jenkins (Server Not Found). Could you please share the error message of failing Jenkins/Precommit check?

From what I see, a few ESIMD and AOT test cases from intel/llvm-test-suite failed at runtime - doesn't seem to be related to this PR, because compilation step was successful and changes in the PR shouldn't affect anything else, they are basically non-functional changes

@bader is correct here: #3111 (comment) - I was looking into the wrong part of the log file

[bader]

Jenkins/Precommit runs llvm-test-suite on accelerators not provided by GitHub CI.

The error is reported for this line: https://github.com/intel/llvm-test-suite/blob/intel/SYCL/Basic/handler/handler_mem_op.cpp#L171

In file included from include/sycl/CL/sycl/ONEAPI/reduction.hpp:14:
include/sycl/CL/sycl/handler.hpp:1825:5: error: static_assert failed due to requirement 'isValidModeForSourceAccessor((sycl::access::mode)1029)' "Invalid source accessor mode for the copy method."
static_assert(isValidModeForSourceAccessor(AccessMode_Src),
^ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
llvm-test-suite/SYCL/Basic/handler/handler_mem_op.cpp:669:11: note: in instantiation of function template specialization 'sycl::handler::copy<float, 1, sycl::access::mode::atomic, sycl::access::target::global_buffer, float, 0, sycl::access::mode::write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t, sycl::access::placeholder::false_t>' requested here
Cgh.copy(AccessorFrom, AccessorTo);
^
llvm-test-suite/SYCL/Basic/handler/handler_mem_op.cpp:171:5: note: in instantiation of function template specialization 'test_0D1D_copy_acc_acc_atomic' requested here
test_0D1D_copy_acc_acc_atomic();
^

It looks like it might be a bug in the reduction implementation. Tagging @v-klochkov and @Pennycook.

[bader]

It looks like it might be a bug in the reduction implementation. Tagging @v-klochkov and @Pennycook.

It looks like SYCL spec doesn't list atomic access mode as allowed in copy methods.

[Pennycook]

It looks like it might be a bug in the reduction implementation. Tagging @v-klochkov and @Pennycook.

It looks like SYCL spec doesn't list atomic access mode as allowed in copy methods.

Agreed. I can't see any issues with the reduction implementation, but the test at https://github.com/intel/llvm-test-suite/blob/intel/SYCL/Basic/handler/handler_mem_op.cpp#L168 is explicitly trying to use atomic as the access mode.

I think this is an issue with the test. I can't see why we would want to allow values to be copied between atomic accessors.

[v-klochkov]

Good catch. Indeed, atomic mode is not listed for the accessor modes that can be passed to handler::copy().
I'll remove the big portion of this patch enabling handler::copy for atomics: #1551

The fix for the LIT test is uploaded: intel/llvm-test-suite#126
I'll update the handler class soon.