File: sycl_atomic.cpp

package info (click to toggle)
taskflow 3.9.0%2Bds-1
  • links: PTS, VCS
  • area: main
  • in suites: forky, sid, trixie
  • size: 45,948 kB
  • sloc: cpp: 39,058; xml: 35,572; python: 12,935; javascript: 1,732; makefile: 59; sh: 16
file content (55 lines) | stat: -rw-r--r-- 1,177 bytes parent folder | download
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
// This program demonstrates how to create a simple vector-add
// application using syclFlow and unified shared memory (USM).

#include <taskflow/sycl/syclflow.hpp>

constexpr size_t N = 10000;

int main() {

  // create a standalone scylFlow
  sycl::queue queue;

  tf::syclFlow syclflow(queue);

  // allocate a shared memory and initialize the data
  auto data = sycl::malloc_shared<int>(N, queue);

  for(size_t i=0; i<N; i++) {
    data[i] = i;
  }

  // reduce the summation to the first element using ONEAPI atomic_ref
  syclflow.parallel_for(
    sycl::range<1>(N), [=](sycl::id<1> id) {

      auto ref = sycl::atomic_ref<
        int, 
        sycl::memory_order_relaxed, 
        sycl::memory_scope::device,
        sycl::access::address_space::global_space
      >{data[0]};

      ref.fetch_add(data[id]);
    }
  );

  // run the syclflow
  syclflow.offload();

  // create a deallocate task that checks the result and frees the memory
  if(data[0] != (N-1)*N/2) {
    std::cout << data[0] << '\n';
    throw std::runtime_error("incorrect result");
  }

  std::cout << "correct result\n";

  // deallocates the memory
  sycl::free(data, queue);


  return 0;
}