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

Most upvoted comments

atomic operators implemented in PR #210