#### Task vs data parallelism
* **Task parallelism** is where you have several,
possibly distinct tasks executing in parallel.
* In task parallelism you optimize for latency.
* **Data parallelism** is where you have the same
task being performed on multiple elements of data.
* In data parallelism you optimize for throughput.
#### Vector processors
* Many processors are vector processors, which means
they can naturally perform data parallelism.
* GPUs are designed to be parallel.
* CPUs have SIMD instructions which perform the
same instruction on a number elements of data.
#### SPMD model for describing data parallelism
Sequential CPU code
void calc(const int in[], int out[]) {
// all iterations are run in the same
// thread in a loop
for (int i = 0; i < 1024; i++){
out[i] = in[i] * in[i];
}
}
// calc is invoked just once and all
// iterations are performed inline
calc(in, out);
Parallel SPMD code
void calc(const int in[], int out[], int id) {
// function is described in terms of
// a single iteration
out[id] = in[id] * in[id];
}
// parallel_for invokes calc multiple
// times in parallel
parallel_for(calc, in, out, 1024);
#### SYCL execution model
* In SYCL kernel functions are executed by
**work-items**.
* You can think of a work-item as a thread of
execution.
* Each work-item will execute a SYCL kernel function from start to end.
* A work-item can run on CPU threads, SIMD lanes,
GPU threads, or any other kind of processing
element.
#### SYCL execution model
* SYCL kernel functions are invoked within an **nd-range**
* An nd-range contains a number of work-groups which are made up of a number of work-items
* Work-groups always have the same number of work-items
#### SYCL execution model
* The nd-range describes an **iteration space** composed in terms of work-groups and work-items
* An nd-range can be 1, 2 or 3 dimensions
* An nd-range has two components
* The **global-range** describes the total number of work-items in each dimension
* The **local-range** describes the number of work-items in a work-group in each dimension
#### SYCL execution model
* Each invocation in the iteration space of an nd-range is a work-item
* Each invocation knows which work-item it is on and can query certain information about its position in the nd-range
* Each work-item has the following:
* **Global range**: {12, 12}
* **Global id**: {5, 6}
* **Group range**: {3, 3}
* **Group id**: {1, 1}
* **Local range**: {4, 4}
* **Local id**: {1, 2}
#### SYCL execution model
Typically an nd-range invocation SYCL will execute the SYCL kernel function on a very large number of work-items, often in the thousands
#### SYCL execution model
* Multiple work-items will generally execute concurrently
* On vector hardware this is often done in lock-step, which means the same hardware instructions
* The number of work-items that will execute concurrently can vary from one device to another
* Work-items will be batched along with other work-items in the same work-group
* The order work-items and work-groups are executed in is implementation defined
#### SYCL execution model
* Work-items in a work-group can be synchronized using a work-group barrier
* All work-items within a work-group must reach the barrier before any can continue on
#### SYCL execution model
* SYCL does not support synchronizing across all work-items in the nd-range
* The only way to do this is to split the computation into separate SYCL kernel functions
#### SYCL execution model
* SYCL also provides a simplified execution model with `sycl::range` in place of `sycl::nd_range`
* Caller only provides the global range
* Local range is decided by the runtime and cannot be inspected
* No synchronization is possible between work items
* Useful for simple problems which don't require synchronization, local memory and ultimate performance
* Runtime may not always have enough information to choose the best-performing size
#### Parallel_for
q.parallel_for(nd_range{{1024, 16}, {32, 4}},
[=](nd_item<2> item){
// SYCL kernel function is executed
// on a range of work-items
});
* In SYCL, kernel functions can be enqueued to execute
over a range of work-items using `parallel_for`
* The first argument to `parallel_for` is an `nd_range` or
a `range` which describes the iteration space over which
the kernel is to be executed
* The kernel function has to take an `nd_item` or `item`,
respectively, as the parameter (or any type they can be
implicitly converted to, commonly from `item` to `id`)
#### Expressing parallelism
q.parallel_for(nd_range<1>{1024,32},
[=](nd_item<1> ndItem){
/* kernel function code */
id globalId = ndItem.get_global_id();
id localId = ndItem.get_local_id();
});
* Overload taking an `nd_range` object specifies the global and local range
* An `nd_item` parameter represents the global and local range and index
q.parallel_for(range<1>{1024},
[=](item<1> item){
/* kernel function code */
id globalId = item.get_id();
});
* Overload taking a `range` object specifies the global range, runtime decides local range
* An `item` parameter represents the global range and the index within the global range
q.parallel_for(range<1>{1024},
[=](id<1> globalId){
/* kernel function code */
});
* Overload taking a `range` object specifies the global range, runtime decides local range
* An `id` parameter represents the index within the global range
#### SYCL memory model
* Each work-item can access a dedicated region of **private memory**
* A work-item cannot access the private memory of another work-item
#### SYCL memory model
* Each work-item can access a dedicated region of **local memory** accessible to all work-items in a work-group
* A work-item cannot access the local memory of another work-group
#### SYCL memory model
* Each work-item can access a single region of **global memory** that's accessible to all work-items in a ND-range
#### SYCL memory model
* Each memory region has a different size and access latency
* Global memory is larger than local memory and local memory is larger than private memory
* Private memory is faster than local memory and local memory is faster than global memory
#### Accessing Data
#### Accessing Data
#include<sycl/sycl.hpp>
#include<vector>
int main() {
std::vector<float> hA{ 3.14 4.38 6.9, 10.2 }, hB{ 6.5, 11.4, 0.2, 9.12 };
queue gpuQueue(gpu_selector{});
float* dA = malloc_shared<float>(hA.size(), gpuQueue);
float* dB = malloc_shared<float>(hB.size(), gpuQueue);
float* out = malloc_shared<float>(hA.size(), gpuQueue);
event cpy_evt_1 = gpuQueue.copy(hA.data(), dA, hA.size());
event cpy_evt_2 = gpuQueue.copy(hB.data(), dB, hB.size());
gpuQueue.parallel_for(range&<1>(dA.size()) {cpy_evt_1, cpy_evt_2},
[=](id<1> i){
out[i] = inA[i] + inB[i];
});
gpuQueue.wait();
/* ... */
}
* The subscript operator can take an **id**
* Here we access the data by
passing in the `id` passed to the SYCL kernel
function into the subscript operator
* `.get_id(dim)` takes a **size_t** to get the index for a single dimension
* E.g. a 3D item: `itm.get_id(0)`, `itm.get_id(1)`, `itm.get_id(2)` for x, y, z
#### Exercise
Code_Exercises/Data_Parallelism/source.cpp
Implement a SYCL application using `parallel_for` to add two arrays of values
* Use USM to manage data
* Try the `sycl::range` and `sycl::nd_range` variants