Data Parallel C++ Mastering DPC++ for Programming of Heterogeneous Systems using C++ and SYCL

Table of Contents
About the Authors
Preface
Acknowledgments
Chapter 1: Introduction
Read the Book, Not the Spec
SYCL 1.2.1 vs. SYCL 2020, and DPC++
Getting a DPC++ Compiler
Book GitHub
Hello, World! and a SYCL Program Dissection
Queues and Actions
It Is All About Parallelism
Throughput
Latency
Think Parallel
Amdahl and Gustafson
Scaling
Heterogeneous Systems
Data-Parallel Programming
Key Attributes of DPC++ and SYCL
Single-Source
Host
Devices
Sharing Devices
Kernel Code
Kernel: Vector Addition (DAXPY)
Asynchronous Task Graphs
Race Conditions When We Make a Mistake
C++ Lambda Functions
Portability and Direct Programming
Concurrency vs. Parallelism
Summary
Chapter 2: Where Code Executes
Single-Source
Host Code
Device Code
Choosing Devices
Method#1: Run on a Device of Any Type
Queues
Binding a Queue to a Device, When Any Device Will Do
Method#2: Using the Host Device for Development and Debugging
Method#3: Using a GPU (or Other Accelerators)
Device Types
Accelerator Devices
Device Selectors
When Device Selection Fails
Method#4: Using Multiple Devices
Method#5: Custom (Very Specific) Device Selection
device_selector Base Class
Mechanisms to Score a Device
Three Paths to Device Code Execution on CPU
Creating Work on a Device
Introducing the Task Graph
Where Is the Device Code?
Actions
Fallback
Summary
Chapter 3: Data Management
Introduction
The Data Management Problem
Device Local vs. Device Remote
Managing Multiple Memories
Explicit Data Movement
Implicit Data Movement
Selecting the Right Strategy
USM, Buffers, and Images
Unified Shared Memory
Accessing Memory Through Pointers
USM and Data Movement
Explicit Data Movement in USM
Implicit Data Movement in USM
Buffers
Creating Buffers
Accessing Buffers
Access Modes
Ordering the Uses of Data
In-order Queues
Out-of-Order (OoO) Queues
Explicit Dependences with Events
Implicit Dependences with Accessors
Choosing a Data Management Strategy
Handler Class: Key Members
Summary
Chapter 4: Expressing Parallelism
Parallelism Within Kernels
Multidimensional Kernels
Loops vs. Kernels
Overview of Language Features
Separating Kernels from Host Code
Different Forms of Parallel Kernels
Basic Data-Parallel Kernels
Understanding Basic Data-Parallel Kernels
Writing Basic Data-Parallel Kernels
Details of Basic Data-Parallel Kernels
The range Class
The id Class
The item Class
Explicit ND-Range Kernels
Understanding Explicit ND-Range Parallel Kernels
Work-Items
Work-Groups
Sub-Groups
Writing Explicit ND-Range Data-Parallel Kernels
Details of Explicit ND-Range Data-Parallel Kernels
The nd_range Class
The nd_item Class
The group Class
The sub_group Class
Hierarchical Parallel Kernels
Understanding Hierarchical Data-Parallel Kernels
Writing Hierarchical Data-Parallel Kernels
Details of Hierarchical Data-Parallel Kernels
The h_item Class
The private_memory Class
Mapping Computation to Work-Items
One-to-One Mapping
Many-to-One Mapping
Choosing a Kernel Form
Summary
Chapter 5: Error Handling
Safety First
Types of Errors
Let’s Create Some Errors!
Synchronous Error
Asynchronous Error
Application Error Handling Strategy
Ignoring Error Handling
Synchronous Error Handling
Asynchronous Error Handling
The Asynchronous Handler
Invocation of the Handler
Errors on a Device
Summary
Chapter 6: Unified Shared Memory
Why Should We Use USM?
Allocation Types
Device Allocations
Host Allocations
Shared Allocations
Allocating Memory
What Do We Need to Know?
Multiple Styles
Allocations à la C
Allocations à la C++
C++ Allocators
Deallocating Memory
Allocation Example
Data Management
Initialization
Data Movement
Explicit
Implicit
Migration
Fine-Grained Control
Queries
Summary
Chapter 7: Buffers
Buffers
Creation
Buffer Properties
use_host_ptr
use_mutex
context_bound
What Can We Do with a Buffer?
Accessors
Accessor Creation
What Can We Do with an Accessor?
Summary
Chapter 8: Scheduling Kernels and Data Movement
What Is Graph Scheduling?
How Graphs Work in DPC++
Command Group Actions
How Command Groups Declare Dependences
Examples
When Are the Parts of a CG Executed?
Data Movement
Explicit
Implicit
Synchronizing with the Host
Summary
Chapter 9: Communication and Synchronization
Work-Groups and Work-Items
Building Blocks for Efficient Communication
Synchronization via Barriers
Work-Group Local Memory
Using Work-Group Barriers and Local Memory
Work-Group Barriers and Local Memory in ND-Range Kernels
Local Accessors
Synchronization Functions
A Full ND-Range Kernel Example
Work-Group Barriers and Local Memory in Hierarchical Kernels
Scopes for Local Memory and Barriers
A Full Hierarchical Kernel Example
Sub-Groups
Synchronization via Sub-Group Barriers
Exchanging Data Within a Sub-Group
A Full Sub-Group ND-Range Kernel Example
Collective Functions
Broadcast
Votes
Shuffles
Loads and Stores
Summary
Chapter 10: Defining Kernels
Why Three Ways to Represent a Kernel?
Kernels As Lambda Expressions
Elements of a Kernel Lambda Expression
Naming Kernel Lambda Expressions
Kernels As Named Function Objects
Elements of a Kernel Named Function Object
Interoperability with Other APIs
Interoperability with API-Defined Source Languages
Interoperability with API-Defined Kernel Objects
Kernels in Program Objects
Summary
Chapter 11: Vectors
How to Think About Vectors
Vector Types
Vector Interface
Load and Store Member Functions
Swizzle Operations
Vector Execution Within a Parallel Kernel
Vector Parallelism
Summary
Chapter 12: Device Information
Refining Kernel Code to Be More Prescriptive
How to Enumerate Devices and Capabilities
Custom Device Selector
Being Curious: get_info<>
Being More Curious: Detailed Enumeration Code
Inquisitive: get_info<>
Device Information Descriptors
Device-Specific Kernel Information Descriptors
The Specifics: Those of “Correctness”
Device Queries
Kernel Queries
The Specifics: Those of “Tuning/Optimization”
Device Queries
Kernel Queries
Runtime vs. Compile-Time Properties
Summary
Chapter 13: Practical Tips
Getting a DPC++ Compiler and Code Samples
Online Forum and Documentation
Platform Model
Multiarchitecture Binaries
Compilation Model
Adding SYCL to Existing C++ Programs
Debugging
Debugging Kernel Code
Debugging Runtime Failures
Initializing Data and Accessing Kernel Outputs
Multiple Translation Units
Performance Implications of Multiple Translation Units
When Anonymous Lambdas Need Names
Migrating from CUDA to SYCL
Summary
Chapter 14: Common Parallel Patterns
Understanding the Patterns
Map
Stencil
Reduction
Scan
Pack and Unpack
Pack
Unpack
Using Built-In Functions and Libraries
The DPC++ Reduction Library
The reduction Class
The reducer Class
User-Defined Reductions
oneAPI DPC++ Library
Group Functions
Direct Programming
Map
Stencil
Reduction
Scan
Pack and Unpack
Pack
Unpack
Summary
For More Information
Chapter 15: Programming for GPUs
Performance Caveats
How GPUs Work
GPU Building Blocks
Simpler Processors (but More of Them)
Expressing Parallelism
Expressing More Parallelism
Simplified Control Logic (SIMD Instructions)
Predication and Masking
SIMD Efficiency
SIMD Efficiency and Groups of Items
Switching Work to Hide Latency
Offloading Kernels to GPUs
SYCL Runtime Library
GPU Software Drivers
GPU Hardware
Beware the Cost of Offloading!
Transfers to and from Device Memory
GPU Kernel Best Practices
Accessing Global Memory
Accessing Work-Group Local Memory
Avoiding Local Memory Entirely with Sub-Groups
Optimizing Computation Using Small Data Types
Optimizing Math Functions
Specialized Functions and Extensions
Summary
For More Information
Chapter 16: Programming for CPUs
Performance Caveats
The Basics of a General-Purpose CPU
The Basics of SIMD Hardware
Exploiting Thread-Level Parallelism
Thread Affinity Insight
Be Mindful of First Touch to Memory
SIMD Vectorization on CPU
Ensure SIMD Execution Legality
SIMD Masking and Cost
Avoid Array-of-Struct for SIMD Efficiency
Data Type Impact on SIMD Efficiency
SIMD Execution Using single_task
Summary
Chapter 17: Programming for FPGAs
Performance Caveats
How to Think About FPGAs
Pipeline Parallelism
Kernels Consume Chip “Area”
When to Use an FPGA
Lots and Lots of Work
Custom Operations or Operation Widths
Scalar Data Flow
Low Latency and Rich Connectivity
Customized Memory Systems
Running on an FPGA
Compile Times
The FPGA Emulator
FPGA Hardware Compilation Occurs “Ahead-of-Time”
Writing Kernels for FPGAs
Exposing Parallelism
Keeping the Pipeline Busy Using ND-Ranges
Pipelines Do Not Mind Data Dependences!
Spatial Pipeline Implementation of a Loop
Loop Initiation Interval
Pipes
Blocking and Non-blocking Pipe Accesses
For More Information on Pipes
Custom Memory Systems
Some Closing Topics
FPGA Building Blocks
Clock Frequency
Summary
Chapter 18: Libraries
Built-In Functions
Use the sycl:: Prefix with Built-In Functions
DPC++ Library
Standard C++ APIs in DPC++
DPC++ Parallel STL
DPC++ Execution Policy
FPGA Execution Policy
Using DPC++ Parallel STL
Using Parallel STL with USM
Error Handling with DPC++ Execution Policies
Summary
Chapter 19: Memory Model and Atomics
What Is in a Memory Model?
Data Races and Synchronization
Barriers and Fences
Atomic Operations
Memory Ordering
The Memory Model
The memory_order Enumeration Class
The memory_scope Enumeration Class
Querying Device Capabilities
Barriers and Fences
Atomic Operations in DPC++
The atomic Class
The atomic_ref Class
Using Atomics with Buffers
Using Atomics with Unified Shared Memory
Using Atomics in Real Life
Computing a Histogram
Implementing Device-Wide Synchronization
Summary
For More Information
Epilogue: Future Direction of DPC++
Alignment with C++20 and C++23
Address Spaces
Extension and Specialization Mechanism
Hierarchical Parallelism
Summary
For More Information
Index