Lesson_Materials/Lesson-1-Introduction-to-SYCL/index.html

<!DOCTYPE html>

<html>
	<head>
	    <meta charset="utf-8">
		<link rel="stylesheet" href="../common-revealjs/css/reveal.css">
		<link rel="stylesheet" href="../common-revealjs/css/theme/white.css">
		<link rel="stylesheet" href="../common-revealjs/css/custom.css">
		<script>
			// This is needed when printing the slides to pdf
			var link = document.createElement( 'link' );
			link.rel = 'stylesheet';
			link.type = 'text/css';
			link.href = window.location.search.match( /print-pdf/gi ) ? 'css/print/pdf.css' : 'css/print/paper.css';
			document.getElementsByTagName( 'head' )[0].appendChild( link );
		</script>
		<script>
		    // This is used to display the static images on each slide,
			// See global-images in this html file and custom.css
			(function() {
				if(window.addEventListener) {
					window.addEventListener('load', () => {
						let slides = document.getElementsByClassName("slide-background");

						if (slides.length === 0) {
							slides = document.getElementsByClassName("pdf-page")
						}

						// Insert global images on each slide
						for(let i = 0, max = slides.length; i < max; i++) {
							let cln = document.getElementById("global-images").cloneNode(true);
							cln.removeAttribute("id");
							slides[i].appendChild(cln);
						}

						// Remove top level global images
						let elem = document.getElementById("global-images");
						elem.parentElement.removeChild(elem);
					}, false);
				}
			})();
		</script>
		
	</head>
	<body>
		<div class="reveal">
			<div class="slides">
				<div id="global-images" class="global-images">
					<img src="../common-revealjs/images/sycl_academy.png" />
					<img src="../common-revealjs/images/sycl_logo.png" />
					<img src="../common-revealjs/images/trademarks.png" />
					<img src="../common-revealjs/images/codeplay.png" />
				</div>
				<!--Slide 1-->
				<section class="hbox">
					<div class="hbox" data-markdown>
						## Introduction to SYCL
					</div>
				</section>
				<!--Slide 2-->
				<section class="hbox" data-markdown>
					## Learning Objectives
					* Learn about the SYCL 1.2.1 specification and it's implementations
					* Learn about the major features that SYCL provides
					* Learn about the components of a SYCL implementation
					* Learn about the anatomy of a typical SYCL application
					* Learn where to find useful resources for SYCL
                </section>
				<!--Slide 3-->
				<section class="hbox" data-markdown>
				    ## What is SYCL?
				</section>
				<!--Slide 4-->
				<section>
					<div class="hbox" data-markdown>
						SYCL is a single source, high-level, standard C++ programming model, that can target a range of heterogeneous platforms
					</div>
					<aside class="notes">
						We'll examine what each of the parts of this statement means.
					</aside>
				</section>
				<!--Slide 5-->
				<section class="hbox">
					<div class="hbox" data-markdown>
						SYCL is a **single source**, high-level, standard C++ programming model, that can target a range of heterogeneous platforms
					</div>
					<div class="container">
						<div class="col" data-markdown>
							![SYCL](../common-revealjs/images/sycl.png "SYCL")
						</div>
						<div class="col" data-markdown>
							*  SYCL allows you to write both host CPU and device code in the same C++ source file
							*  This requires two compilation passes; one for the host code and one for the device code
						</div>
					</div>
					<aside class="notes">
						Single source means that SYCL kernel source code can live in the same source file as the other code you are using.
						This is not the case with something like OpenCL, another heterogeneous programming interface, where the kernel code lives in a separate file.
						Your kernel contains the code you want to execute on parallel hardware, i.e the bit you want to accelerate such as a vector addition
					</aside>
				</section>
				<!--Slide 6-->
				<section class="hbox">
					<aside class="notes">
						TBD
					</aside>
					<div class="hbox" data-markdown>
						SYCL is a single source, **high-level**, standard C++ programming model, that can target a range of heterogeneous platforms
					</div>
					<div class="container">
					<aside class="aside aside-1" data-markdown>
						![SYCL](../common-revealjs/images/sycl_opencl.png "SYCL-OpenCL")
					</aside>
					<article class="main" data-markdown>
						*  SYCL provides high-level abstractions over common boilerplate code 
						  * Platform/device selection
						  * Buffer creation
						  * Kernel compilation
						  * Dependency management and scheduling
					</article>
					<aside class="notes">
						
					</aside>
				</section>
				<!--Slide 7-->
				<section class="hbox">
				    <div class="hbox" data-markdown>
				        SYCL is a single source, high-level **standard C++** programming model, that can target a range of heterogeneous platforms
					</div>
					<div class="container">
						<div class="col" data-markdown>
							![SYCL](../common-revealjs/images/code_comparison.png "SYCL-Comparison")
						</div>
						<div class="col" data-markdown>
							* SYCL allows you to write standard C++
							* Unlike the other implementations shown on the left there are:
							  * No language extensions 
							  * No pragmas 
							  * No attributes
						</div>
					</div>
					<aside class="notes">
						SYCL only uses standard C++ code to contruct the code used to dispatch your work
						As you can see other models use proprietary things like pragmas, attributes and language extensions
					</aside>
				</section>
				<!--Slide 8-->
				<section class="hbox">
					<div class="hbox" data-markdown>
						SYCL is a single source, high-level standard C++ programming model, that can **target a range of heterogeneous platforms**
					</div>
					<div class="container">
						<div class="col" data-markdown>
							![SYCL](../common-revealjs/images/sycl_targets.png "SYCL-Targets")
						</div>
						<div class="col" data-markdown>
							* SYCL can target any device supported by its back-end
							* SYCL can target a number of different backends
							
							While the current specification is limited to OpenCL, some implementations are already supporting other non-OpenCL back-ends.
						</div>
					</div>
					<aside class="notes">
						SYCL has been designed to enable developers to write C++ code that can be used to target different types of processors that are part of a heterogeneous system. 
						This means that you can write your code and deploy it to devices that support the SYCL implementation's back-end. 
						Whilst the current specification of SYCL is bound to OpenCL, it is possible to support non-OpenCL back-ends. For example hipSYCL targets AMD's HIP interface and ComputeCpp targets NVidia's ptx instruction set
					</aside>
				</section>
				<!--Slide 9-->
				<section data-markdown>
					#### What SYCL implementations are available?
					![SYCL Implementations](https://raw.githubusercontent.com/illuhad/hipSYCL/master/doc/img/sycl-targets.png "SYCL-Implementations")
					Image referenced from the hipSYCL project (https://github.com/illuhad/hipSYCL)
				</section>
				<!--Slide 10-->
				<section data-markdown>
					## What is in a SYCL Implementation?
				</section>
				<!--Slide 11-->
				<section class="hbox">
					<div class="container">
						<div class="col" data-markdown>
							![SYCL Implementation](../common-revealjs/images/sycl_implementation_sycl_interface.png "SYCL-Implementation")
						</div>
						<div class="col" data-markdown>
							* The SYCL interface is a C++ template library that developers can use to access the features of SYCL
							* The same interface is used for both the host and device code
						</div>
					</div>
					<div class="bottom-bullets" data-markdown>
						* The host is generally the CPU and is used to dispatch the parallel execution of kernels
						* The device is the parallel unit used to execute the kernels, such as a GPU
					</div>
				</section>
				<!--Slide 12-->
				<section class="hbox">
					<div class="container">
						<div class="col" data-markdown>
							![SYCL Runtime](../common-revealjs/images/sycl_implementation_sycl_runtime.png "SYCL-Host-Device")
						</div>
						<div class="col" data-markdown>
							* The SYCL runtime is a library that schedules and executes work
							  * It loads kernels, tracks data dependencies and schedules commands
						</div>
				</section>
				<!--Slide 13-->
				<section class="hbox">
					<div class="container">
						<div class="col" data-markdown>
							![SYCL Host Device](../common-revealjs/images/sycl_implementation_host_device.png "SYCL-Host-Device")
						</div>
						<div class="col" data-markdown>
							* The host device is an emulated backend that is executed as native C++ code and emulates the SYCL execution and memory model 
							* The host device can be used to execute kernels without backend drivers and for debugging purposes
						</div>
				</section>
				<!--Slide 14-->
				<section class="hbox">
					<div class="container">
						<div class="col" data-markdown>
							![SYCL Backend](../common-revealjs/images/sycl_implementation_backend_interface.png "SYCL-Backend")
						</div>
						<div class="col" data-markdown>
							* The back-end interface is where the SYCL runtime calls down into a back-end in order to execute on a particular device  
							* The standard back-end is OpenCL but some implementations support other interfaces
						</div>
					</div>
				</section>
				<!--Slide 15-->
				<section class="hbox">
					<div class="container">
						<div class="col" data-markdown>
							![SYCL Compiler](../common-revealjs/images/sycl_implementation_sycl_compiler.png "SYCL-Compiler")
						</div>
						<div class="col" data-markdown>
							* The SYCL device compiler is a C++ compiler which can identify SYCL kernels and compile them down to an IR or ISA
							  * This can be SPIR, SPIR-V, GCN, PTX or any proprietary vendor ISA
						</div>
					</div>
					<div class="bottom-bullets" data-markdown>
						**IR** = Intermediate Representation 	**ISA** = Instruction Set Architecture
					</div>
				</section>
				<!--Slide 16-->
				<section data-markdown>
					## What does a SYCL Application Look Like?
				</section>
				<!--Slide 17-->
				<section>
					
					<div class="hbox" data-markdown>
						#### Include the SYCL Header File
					</div>
					<div class="hbox">
						<code><pre>

<mark>#include &ltCL/sycl.hpp&gt </mark>
using namespace cl::sycl;

int main(int argc, char *argv[]) { 

}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
							* First we include the SYCL header which contains the runtime API
							* We also import the **cl::sycl** namespace here, this reduces the amount of code we need to write
					</div>
				</section>
				<!--Slide 18-->
				<section>
					<div class="hbox" data-markdown>
						#### Device Selectors And Queues
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt
using namespace cl::sycl;

int main(int argc, char *argv[]) { 

  <mark>queue gpuQueue{gpu_selector{}};</mark>

}
						</code></pre>
					</div>
					<div class="container">
						<div class="col-left" data-markdown>
							* Device selectors allow you to choose a device based on a custom configuration
							* The queue default constructor uses a the default_selector, which allows the runtime to select a device for you
						</div>
						<div class="col-right" data-markdown>
							![SYCL Selector](../common-revealjs/images/sycl_code_selector.png "SYCL-Queues")
						</div>
					</div>
				</section>
				<!--Slide 19-->
				<section>
					<div class="hbox" data-markdown>
						#### Queues and Command Groups
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt
using namespace cl::sycl; 

int main(int argc, char *argv[]) { 

  queue gpuQueue{gpu_selector{}};

  <mark>gpuQueue.submit([&](handler &cgh){ </mark>

  });
}
						</code></pre>
					</div>
					<div class="container">
						<div class="col-left" data-markdown>
							With a queue we can submit a command group; a command group contains
							* A SYCL command (e.g. a SYCL kernel function) 
							* Execution range 
							* Accessors
						</div>
						<div class="col-right" data-markdown>
							![SYCL Selector](../common-revealjs/images/sycl_code_cg.png "SYCL-Queues")
						</div>
					</div>
				</section>
				<!--Slide 20-->
				<section>
					<div class="hbox" data-markdown>
						#### Set up Vectors
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt 
using namespace cl::sycl; 

int main(int argc, char *argv[]) { 
  <mark>std::vector&ltfloat&gt dA{ … }, dB{ … }, dO{ … };</mark>

  queue gpuQueue{gpu_selector{}};

  gpuQueue.submit([&](handler &cgh){

  });
}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
						We initialize three vectors, two inputs (**dA**, **dB**) and an output (**dO**)
					</div>
				</section>
				<!--Slide 21-->
				<section>
					<div class="hbox" data-markdown>
						#### Create Buffers
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt 
using namespace cl::sycl; 

int main(int argc, char *argv[]) { 
  std::vector&ltfloat&gt dA{ … }, dB{ … }, dO{ … }; // psuedo code

  queue gpuQueue{gpu_selector{}};

  <mark>buffer&ltfloat, 1&gt bufA(dA.data(), range&lt1&gt(dA.size()));</mark>
  <mark>buffer&ltfloat, 1&gt bufB(dB.data(), range&lt1&gt(dB.size()));</mark>
  <mark>buffer&ltfloat, 1&gt bufO(dO.data(), range&lt1&gt(dO.size()));</mark>

  gpuQueue.submit([&](handler &cgh){

  });
}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
					    We create a buffer for each vector to manage the data across host and device
					</div>
				</section>
				<!--Slide 22-->
				<section>
					<div class="hbox" data-markdown>
						#### Buffers on Leaving Scope
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt
using namespace cl::sycl; 

int main(int argc, char *argv[]) { 
  std::vector<float> dA{ … }, dB{ … }, dO{ … };

  queue gpuQueue{gpu_selector{}};

  <mark>{</mark>
    buffer<float, 1> bufA(dA.data(), range<1>(dA.size())); 
    buffer<float, 1> bufB(dB.data(), range<1>(dB.size())); 
    buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));

    gpuQueue.submit([&](handler &cgh){

    });
  <mark>}</mark>
}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
					    Buffers synchronize on destruction via RAII
						So adding this scope means that all kernels writing to the buffers will wait and the data will be copied back to the vectors on leaving this scope
					</div>
				</section>
				<!--Slide 23-->
				<section>
					<div class="hbox" data-markdown>
						#### Create Accessors
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt
using namespace cl::sycl;

int main(int argc, char *argv[]) { 
  std::vector<float> dA{ … }, dB{ … }, dO{ … };

  queue gpuQueue{gpu_selector{}};

  {
    buffer<float, 1> bufA(dA.data(), range<1>(dA.size())); 
    buffer<float, 1> bufB(dB.data(), range<1>(dB.size())); 
    buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));

    gpuQueue.submit([&](handler &cgh){

      <mark>auto inA = bufA.get_access&ltaccess::mode::read&gt(cgh);</mark>
      <mark>auto inB = bufB.get_access&ltaccess::mode::read&gt(cgh);</mark>
      <mark>auto out = bufO.get_access&ltaccess::mode::write&gt(cgh);</mark>

	});
  }
}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
						We create an accessor for each of the buffers
						Read access for the two input buffers and write access for the output buffer
					</div>
				</section>
				<!--Slide 24-->
				<section>
					<div class="hbox" data-markdown>
						#### SYCL Kernel
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt
using namespace cl::sycl;
<mark>class add;</mark>

int main(int argc, char *argv[]) { 
  std::vector&ltfloat&gt dA{ … }, dB{ … }, dO{ … };

  queue gpuQueue{gpu_selector{}};

  {
    buffer&ltfloat, 1&gt bufA(dA.data(), range&lt1&gt(dA.size()));
    buffer&ltfloat, 1&gt bufB(dB.data(), range&lt1&gt(dB.size()));
    buffer&ltfloat, 1&gt bufO(dO.data(), range&lt1&gt(dO.size()));

    gpuQueue.submit([&](handler &cgh){

      auto inA = bufA.get_access&ltaccess::mode::read&gt(cgh);
      auto inB = bufB.get_access&ltaccess::mode::read&gt(cgh);
      auto out = bufO.get_access&ltaccess::mode::write&gt(cgh);

      <mark>cgh.parallel_for&ltadd&gt(range&lt1&gt(dA.size()),</mark>
        <mark>[=](id&lt1&gt i){ out[i] = inA[i] + inB[i]; });</mark>

	});
  }
}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
						* We define a SYCL kernel function for the command group using the parallel_for API
						* The first argument here is a range, specifying the iteration space
						* The second argument is a lambda function that represents the entry point for the SYCL kernel
						* This lambda takes an id parameter that describes the current iteration being executed
					</div>
				</section>
				<!--Slide 25-->
				<section>
					<div class="hbox" data-markdown>
						#### Template Parameter
					</div>
					<div class="container">
						<code><pre>
#include &ltCL/sycl.hpp&gt
using namespace cl::sycl;
<mark>class add;</mark>

int main(int argc, char *argv[]) { 
  std::vector&ltfloat&gt dA{ … }, dB{ … }, dO{ … };

  queue gpuQueue{gpu_selector{}};

  {
    buffer&ltfloat, 1&gt bufA(dA.data(), range&lt1&gt(dA.size()));
    buffer&ltfloat, 1&gt bufB(dB.data(), range&lt1&gt(dB.size()));
    buffer&ltfloat, 1&gt bufO(dO.data(), range&lt1&gt(dO.size()));

    gpuQueue.submit([&](handler &cgh){

      auto inA = bufA.get_access<access::mode::read>(cgh);
      auto inB = bufB.get_access<access::mode::read>(cgh);
      auto out = bufO.get_access<access::mode::write>(cgh);

      cgh.parallel_for<mark>&ltadd&gt</mark>(range&lt1&gt(dA.size()),
        [=](id&lt1&gt i){ out[i] = inA[i] + inB[i]; });

    });
  }
}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
						* The template parameter to parallel_for is used to name the lambda
						* The reason for this is that C++ does not have a standard ABI for lambdas so they are represented differently across the host and device compiler
						* SYCL kernel functions follow C++ ODR rules, which means that if a SYCL kernel is in a template context, the kernel name needs to reflect that context, so must contain the same template arguments
					</div>
				</section>
				<!--Slide 26-->
				<section>
					<div class="hbox" data-markdown>
						#### Error Handling
					</div>
					<div class="container">
						<code class="code-60pc"><pre>
#include &ltCL/sycl.hpp&gt
using namespace cl::sycl;
class add;

int main(int argc, char *argv[]) { 
  std::vector&ltfloat&gt dA{ … }, dB{ … }, dO{ … };

  <mark>try{</mark>
    queue gpuQueue{gpu_selector{}, <mark>async_handler{}</mark>};

	{
      buffer&ltfloat, 1&gt bufA(dA.data(), range&lt1&gt(dA.size()));
      buffer&ltfloat, 1&gt bufB(dB.data(), range&lt1&gt(dB.size()));
      buffer&ltfloat, 1&gt bufO(dO.data(), range&lt1&gt(dO.size()));

      gpuQueue.submit([&](handler &cgh){

        auto inA = bufA.get_access<access::mode::read>(cgh);
        auto inB = bufB.get_access<access::mode::read>(cgh);
        auto out = bufO.get_access<access::mode::write>(cgh);

        cgh.parallel_for&ltadd&gt(range&lt1&gt(dA.size()),
          [=](id&lt1&gt i){ out[i] = inA[i] + inB[i]; });

      });
      <mark>gpuQueue.wait_and_throw();</mark>
    <mark>} catch (...) { /* handle errors */ }</mark>
  }
}
						</code></pre>
					</div>
					<div class="bottom-bullets" data-markdown>
					    * In SYCL errors are handled using exception handling, so you should always wrap SYCL code in a try-catch block
						* Some exceptions are thrown synchronously at the point of using a SYCL API
						* Other exceptions are asynchronous and are stored by the runtime and passed to an **async handler** when the queue is told to throw
					</div>
				</section>
				<!--Slide 27-->
				<section data-markdown>
					## Where to Get Started with SYCL
					
					* Visit https://www.khronos.org/sycl/ to find the latest SYCL specifications
					* Checkout the documentation provided with one of the SYCL implementations.
					* Visit https://sycl.tech to find out about all the SYCL implementations, news and videos
				</section>
				<!--Slide 28-->
				<section data-markdown>
					## Questions
				</section>
			</div>
		</div>
		<script src="../common-revealjs/js/reveal.js"></script>
		<script src="../common-revealjs/plugin/markdown/marked.js"></script>
		<script src="../common-revealjs/plugin/markdown/markdown.js"></script>
		<script src="../common-revealjs/plugin/notes/notes.js"></script>
		<script>
			Reveal.initialize({mouseWheel: true, defaultNotes: true});
			Reveal.configure({ slideNumber: true });
		</script>
	</body>
</html>