Sycl Academy LogoSycl Academy Logo
listShow Lessons List
  1. What is SYCL
  2. Enqueueing a Kernel
  3. Managing Data
  4. Handling Errors
  5. Device Discovery
  6. Data Parallelism
  7. Introduction to USM
  8. Using USM
  9. Asynchronous Execution
  10. Data and Dependencies
  11. In Order Queue
  12. Advanced Data Flow
  13. Multiple Devices
  14. ND Range Kernels
  15. Image Convolution
  16. Coalesced Global Memory
  17. Vectors
  18. Local Memory Tiling
  19. Further Optimisations
  20. Matrix Transpose
  21. More SYCL Features
  22. Functors
Lesson
GitHub Iconedit

More SYCL Features

This section shows the use of some SYCL features in order to perform a vector reduce operation.

reduce_naive.cpp shows a reduction performed without using atomics or any other SYCL feature. It uses two separate kernels to achieve device synchronization between reductions step 1 and 2.

SYCL Atomics

reduce_atomic.cpp shows a reduction operation using local memory and global memory atomics. Workgroups perform a local reduction using local atomics and the result is added together using a single global memory sycl::atomic_ref.

SYCL Group Algorithms

reduce_group_algorithm.cpp replaces the local memory atomic add with a call to sycl::reduce_over_group, which can be expected to perform a similar operation to the device function workgroup_reduce_local_mem as seen in reduce_naive.cpp.

SYCL Reductions

reduce_sycl_reduction.cpp replaces the entire logic of the reduction with a SYCL reducer object. The reduction is a parameter of the parallel_for invocation, and there is another reducer object used as a parameter to the kernel lambda.