local_accessors.cpp

#include <CL/sycl.hpp>
#include <iostream>

using mode = sycl::access_mode;
using sycl::access::fence_space;
using sycl::access::target;

static constexpr int blockSize = 32;
static constexpr int numBlocks = 8;

template <typename T>
sycl::local_ptr<T>
getLocalPtr(const sycl::accessor<T, 1, mode::read_write, target::local> &acc) {
  return acc.get_pointer();
}

template <typename T> cl::sycl::local_ptr<T> getLocalPtr(std::nullptr_t) {
  return nullptr;
}

template <bool useLocalMem>
auto myKernel(sycl::handler &cgh, sycl::buffer<int, 1> &buf) {
  auto gm_data = buf.get_access<mode::write>(cgh);
  auto acc_sm_buf = [&]() {
    if constexpr (useLocalMem)
      return sycl::accessor<float, 1, mode::read_write, target::local>{
          sycl::range<1>{blockSize}, cgh};
    else
      return nullptr;
  }();
  return [=](sycl::nd_item<1> itemIdx) {
    sycl::local_ptr<float> sm_buf = getLocalPtr<float>(acc_sm_buf);
    int i = itemIdx.get_local_linear_id();
    if constexpr (useLocalMem) {
      sm_buf[i] = i;
      itemIdx.barrier(fence_space::local_space);
      i = sm_buf[(i + 1) % blockSize];
    } else {
      i = (i + 1) % blockSize;
    }
    gm_data[itemIdx.get_global_linear_id()] = i;
  };
}

template <bool> class Kernel;

template <bool useLocalMem>
void runKernel(sycl::queue &q, sycl::buffer<int, 1> &buffer) {
  const sycl::nd_range<1> range{{numBlocks * blockSize}, {blockSize}};
  q.submit([&](sycl::handler &cgh) {
     cgh.parallel_for<Kernel<useLocalMem>>(range,
                                           myKernel<useLocalMem>(cgh, buffer));
   }).wait_and_throw();
}

int main() {
  sycl::device dev{sycl::gpu_selector{}};

  std::cout << "Device name: " << dev.get_info<sycl::info::device::name>()
            << std::endl;
  sycl::queue q{dev};
  sycl::buffer<int, 1> buffer(numBlocks * blockSize);

  std::cout << "Running kernel with SLM: " << std::endl;
  runKernel<true>(q, buffer);
  std::cout << "    Done" << std::endl;

  std::cout << "Running kernel without SLM: " << std::endl;
  runKernel<false>(q, buffer);
  std::cout << "    Done" << std::endl;

  return 0;
}