* In SYCL the USM (unified shared memory) model can be used to manage data. * Which model you choose can have an effect on how you enqueue kernel functions.

* A GPU has its own memory, separate to CPU memory. * In order for the GPU to use memory from the CPU, the following actions must take place (either explicitly or implicitly): * Memory allocation on the GPU. * Data migration from the CPU to the allocation on the GPU. * Some computation on the GPU. * Migration of the result back to the CPU.

* Memory transfers between CPU and GPU are a bottleneck. * We want to minimize these transfers, when possible.

* There are different ways USM memory can be allocated: host, device and shared.

| Type | Description | Access host | Access device | Nominal location | |--------|---------------------------------|----------------|----------------|----------------------| | device | device global allocations | ✗ | ✓ | device | | host | host allocations | ✓ | ✓ | host | | shared | Allocations shared between both | ✓ | ✓ | migrates as required |

// Allocate memory on device
T *device_ptr = sycl::malloc_device<T>(n, myQueue);

// Copy data to device
myQueue.memcpy(device_ptr, cpu_ptr, n * sizeof(T));

// ...
// Do some computation on device
// ...

// Copy data back to CPU
myQueue.memcpy(result_ptr, device_ptr, n * sizeof(T)).wait();

// Free allocated data
sycl::free(device_ptr, myQueue);
						

* It is important to free memory after it has been used to avoid memory leaks.

// Allocate shared memory
T *shared_ptr = sycl::malloc_shared<T>(n, myQueue);

// Shared memory can be accessed on host as well as device
for (auto i = 0; i < n; ++i)
  shared_ptr[i] = i;

// ...
// Do some computation on device
// ...

// Free allocated data
sycl::free(shared_ptr, myQueue);
						

* Shared memory is accessible on host and device. * Performance of shared memory accesses may be poor depending on platform.

gpuQueue.submit([&](handler &cgh){
  cgh.single_task<mykernel>([=]{
    out[0] = inA[0] + inB[0];
  });
});
						

* Data can be accessed from within a kernel function by calling `operator[]`. * `operator[]` for USM pointers must take a `size_t`.

Code_Exercises/Managing_Data/source

Implement a SYCL application that adds two variables and returns the result using the USM memory model.