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 git@github.com:intel/llvm.git 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/projectdirs/nstaff/cookbg/pe/shasta-21.11/modulefiles/core
module avail llvm/nightly

Vector Addition Example

main.cpp

#include <CL/sycl.hpp>

#include <cmath>
#include <iostream>

namespace sycl = cl::sycl;

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};

  {
    auto a = b_a.get_access<sycl::access::mode::discard_write>();
    auto b = b_b.get_access<sycl::access::mode::discard_write>();    
    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) {
      auto a = b_a.get_access<sycl::access::mode::read>(h);
      auto b = b_b.get_access<sycl::access::mode::read>(h);
      auto c = b_c.get_access<sycl::access::mode::write>(h);

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

  {
    double sum = 0.0;
    auto c = b_c.get_access<sycl::access::mode::read>();
    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 0;
    } else {
      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 purge
$ module use /global/cfs/projectdirs/nstaff/cookbg/pe/shasta-21.11/modulefiles/core
$ module load llvm/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