Description
Hello!
This is not a bug, but a case of a very unclear error message from the SYCL runtime in case of incorrect code.
Example code
Please note, that the code is invalid since a placeholder accessor is not bound using cgh.require before accessing it in a kernel.
While such a mistake is not hard to notice here, it might easily creep into complicated kernels with multiple accessors.
#include <CL/sycl.hpp>
using cl::sycl::access::mode;
using cl::sycl::access::placeholder;
using cl::sycl::access::target;
template <class T, enum mode Mode>
using PlaceholderAccessor =
cl::sycl::accessor<T, 1, Mode, target::global_buffer, placeholder::true_t>;
auto myKernel(cl::sycl::handler &cgh,
PlaceholderAccessor<int, mode::write> data) {
// Oops, forgot to call cgh.require(data) here
return [=](cl::sycl::item<1> itemIdx) {
const int v = itemIdx.get_linear_id();
data[v] = v;
};
}
class KernelName;
int main() {
cl::sycl::device dev{cl::sycl::gpu_selector()};
cl::sycl::queue q{dev};
cl::sycl::buffer<int, 1> buffer(1);
cl::sycl::event e = q.submit([&](cl::sycl::handler &cgh) {
auto kernel = myKernel(cgh, buffer);
cgh.parallel_for<KernelName>(1, kernel);
});
e.wait_and_throw();
return 0;
}Current output
The error thrown when running the code above (using d1d6c5d version of intel/llvm) is rather cryptic (#1927 is mildly related), and requires digging through the code of the runtime:
$ clang++ test.cpp -fsycl -o test && ./test
terminate called after throwing an instance of 'cl::sycl::runtime_error'
what(): Alloca for command not found -59 (CL_INVALID_OPERATION)
Aborted (core dumped)
Expectations
Neither SYCL 1.2.1 nor 2020-provisional specs are explicit about what happens when using a placeholder accessor in the kernel without requiring it. So it seems to fall into undefined/unspecified territory.
However, the current behavior leaves a lot to be desired. A more explicit run-time error, at least suggesting that the issue is related to accessors, would improve things drastically.
Something like try-catch-rethrow here might be a quick fix that will already help a lot.
Ideally, one would prefer to get a compile-time warning in such clear-cut cases as the code above.