AdaptiveCpp: unsupported atomic accessors ?
This code snippet:
// construct a SYCL queue for the selected device
auto queue = cl::sycl::queue(device);
// SYCL buffers
int counter = 0;
auto counter_buf = cl::sycl::buffer<int>(& counter, 1);
// submit the kernel to the queue
queue.submit([&](cl::sycl::handler &cgh) {
// access the SYCL buffers from the device kernel
auto counter_d = counter_buf.get_access<cl::sycl::access::mode::atomic>(cgh);
// launch the kernel
cgh.parallel_for_work_group<class count_groups>(
cl::sycl::range<1>{num_work_groups},
cl::sycl::range<1>{work_group_size},
[=](cl::sycl::group<1> group) {
// print the id of all the groups
printf("group id: %lu\n", group.get_id(0));
// print the id of all the threads
group.parallel_for_work_item([&](cl::sycl::h_item<1> item) {
printf("global thread id: %zu\n", item.get_global_id(0));
});
cl::sycl::atomic_fetch_add(counter_d[0], 1);
});
it fails with
test.cc:53:43: error: no viable overloaded operator[] for type 'const cl::sycl::accessor<int, 1, cl::sycl::access::mode::atomic, cl::sycl::access::target::global_buffer,
cl::sycl::access::placeholder::false_t>'
cl::sycl::atomic_fetch_add(counter_d[0], 1);
~~~~~~~~~^~
and in fact the atomic accessors in include/CL/sycl/accessor.hpp are commented out:
/* Available only when: accessMode == access::mode::atomic && dimensions == 0*/
//operator atomic<dataT, access::address_space::global_space> () const;
/* Available only when: accessMode == access::mode::atomic && dimensions > 0*/
//atomic<dataT, access::address_space::global_space> operator[](
// id<dimensions> index) const;
//atomic<dataT, access::address_space::global_space> operator[](
// size_t index) const;
What is the reason for that ?
About this issue
- Original URL
- State: closed
- Created 5 years ago
- Comments: 15
atomic
operators implemented in PR #210