Skip to content

SYCL

SYCL is a cross-platform abstraction layer that enables code for heterogeneous processors to be written using C++ with the host and kernel code for an application contained in the same source file.

Nightly builds

Nightly builds of https://github.com/intel/llvm are available as modules on Perlmutter.

Tip

These modules are also configured to support OpenMP offload on A100!

Warning

A100 support is under active developement!

module use /global/cfs/cdirs/nstaff/cookbg/intel-llvm-sycl-nightly/modulefiles
module avail intel-llvm

Vector Addition Example

main.cpp

#include <CL/sycl.hpp>

#include <cmath>
#include <iostream>

int main() {
  const int n = 100000;
  const sycl::range<1> m{n};

  sycl::buffer<double, 1> b_a{n}, b_b{n}, b_c{n};

  {
    sycl::host_accessor a{b_a, sycl::write_only};
    sycl::host_accessor b{b_b, sycl::write_only};
    for (size_t i = 0; i < n; i++) {
      a[i] = sin(i) * sin(i);
      b[i] = cos(i) * cos(i);
    }
  }

  sycl::queue q{sycl::gpu_selector{}};

  q.submit([&](sycl::handler &h) {
    sycl::accessor a{b_a, h, sycl::read_only};
    sycl::accessor b{b_b, h, sycl::read_only};
    sycl::accessor c{b_c, h, sycl::write_only};

    h.parallel_for(m, [=](sycl::id<1> i) { c[i] = a[i] + b[i]; });
  });

  {
    double sum = 0.0;
    sycl::host_accessor c{b_c, sycl::read_only};
    for (size_t i = 0; i < n; i++)
      sum += c[i];
    std::cout << "sum = " << sum / n << std::endl;

    if (!(fabs(sum - static_cast<double>(n)) <= 1.0e-8))
      return 1;
  }

  return 0;
}

Makefile

CXX = clang++
CXXFLAGS = -std=c++17 -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend '--cuda-gpu-arch=sm_80'

all: sycl-vecadd-buffer.x

sycl-vecadd-buffer.x: main.cpp
    $(CXX) $(CXXFLAGS) -o $@ $^

.PHONY: test
test:
    ./sycl-vecadd-buffer.x

$ module use /global/cfs/cdirs/nstaff/cookbg/intel-llvm-sycl-nightly/modulefiles
$ module load intel-llvm/sycl-nightly/latest
$ make
clang++ -std=c++17 -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend '--cuda-gpu-arch=sm_80' -o sycl-vecadd-buffer.x main.cpp
$ ./sycl-vecadd-buffer.x
sum = 1

References

Support