The SYCL Specification (Sec. 4.9.1) mandates that:

When constructing multi-dimensional ids or ranges from integers, the elements are written such that the right-most element varies fastest in a linearization of the multi-dimensional space.

For this reason in Intel DPC++, the right-most dimension is mapped to the CUDA or HIP x-dimension, the second right-most dimension maps to the CUDA or HIP y-dimension and so on, for example:

cgh.parallel_for(sycl::nd_range{sycl::range(WG_X), sycl::range(WI_X)}, ...)

cgh.parallel_for(sycl::nd_range<2>{sycl::range<2>(WG_Y, WG_X), sycl::range<2>(WI_Y, WI_X)}, ...)

cgh.parallel_for(sycl::nd_range<3>{sycl::range<3>(WG_Z, WG_Y, WG_X), sycl::range<3>(WI_Z, WI_Y, WI_X)}, ...)

Where WG_X and WI_X are the number of work-groups and work-items per work-group (i.e., named grid size and threads per block in CUDA) in the x-dimension and so on for _Y and _Z.

Be aware that this is particularly important in the case of a parallel_for execution of two or three-dimensional kernels when:

• There are manually linearized accesses in (1-d) arrays within local or global memory. This needs to be considered in order to avoid performance issues due to non-coalesced global memory accesses or bank conflicts in local memory. Further details on linearization can be found in Sec. 3.11, Multi-dimensional objects and linearization of the SYCL Specification.

• The following error (or similar) is present:

  Number of work-groups exceed limit for dimension 1 : 379957 > 65535
  
  

  which relates to the fact that on some platforms such as CUDA, the x-dimension can usually support a higher number of work-groups compared to the y and z ones.

  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  
  

Further details on this topic can be found here.

When building for Nvidia GPUs it is possible to omit the specific GPU architecture, which means the compiler will default to sm_50. When using a newer GPU it may be beneficial to specify the exact architecture, which may allow the compiler to use newer more performant features. Refer to the get started guide for instructions on how to specify a specific GPU architecture.

The degree to which DPC++ automatically inlines functions is chosen to balance performance with compilation time. However programmers can also add the always_inline attribute on specific functions to force their inlining, for example:

__attribute__((always_inline)) void function(...) {
  ...
}

...

q.submit([&](sycl::handler &cgh) {
    cgh.parallel_for(..., [=](...) {
      function(...);
    });
}

The SYCL math builtins are defined to match the precision requirements of the equivalent OpenCL 1.2 math builtins, which may be unnecessarily precise for some applications, causing an avoidable loss of performance.

To address this, the SYCL specification provides a native version of a subset of the math functions (full list in Section 4.17.5, Math functions), which trades precision for performance. They are defined within the native namespace. For example, the native version of sycl::cos() is sycl::native::cos().

In general if precision is not an issue using the native variants may provide significant improvements, although do note that not all backends make use of the relaxed precision for all builtins.

When targeting NVIDIA architectures, the following sycl::native:: functions are implemented using the .approx variants of the corresponding PTX instructions:

• sycl::native::divide

• sycl::native::sqrt

• sycl::native::sin

• sycl::native::cos

• sycl::native::log2

• sycl::ext::oneapi::experimental::native::exp2

• sycl::ext::oneapi::experimental::native::tanh

There is a 1:1 mapping between the above sycl::native:: math functions and the PTX instructions, e.g., a call to sycl::native::exp2 results in the compiler generating a single ex2.approx instruction. In other cases, a native math function can be implemented through multiple .approx PTX instructions, for example, sycl::native::tan() is implemented with sin.approx, cos.approx, and divide.approx. It must be noted that native math functions can be faster than their sycl:: counterparts at the cost of losing precision; further details on the precision of the .approx PTX instructions can be found in the PTX ISA documentation.

While the compiler will handle some loop unrolling automatically, it can sometimes be beneficial to help the compiler by manually tuning the unrolling of the compute intensive loops in the device code, for example by using the unrolling pragma as follows:

#pragma unroll <unroll factor>
for( ... ) {
 ...
}

Alias Analysis can prove that two memory references do not alias each other. This may enable optimizations. By default, the compiler must assume that memory references do alias, if not proven otherwise by the alias analysis. It is also possible, however, to explicitly signal to the compiler that a memory reference inside the device code is not aliased. This can be achieved using respective keywords for the buffer/accessor and USM model.

For the former, one can add the no_alias property from the oneapi extension to an accessor:

q.submit([&](sycl::handler &cgh) {
  sycl::accessor acc{..., sycl::ext::oneapi::accessor_property_list{sycl::ext::oneapi::no_alias}};
  ...
});

For the latter, the __restrict__ qualifier can be added to a pointer.

Note that __restrict__ is non-standard C++ and may not behave consistently across sycl implementations. For dpc++ only restrict-qualified device function (a function called from within a SYCL kernel) parameters will be taken into account.

For example:

void function(int *__restrict__ ptr) {
    ...
}

...
int *ptr = sycl::malloc_device<int>(..., q);
...
q.submit([&](sycl::handler &cgh) {
      cgh.parallel_for(..., [=](...) {
        function(ptr);
      });
});

A more brute force approach is to add the [[intel::kernel_args_restrict]] attribute to a kernel. This signals the compiler to ignore all possible alias dependencies between each of the USM pointers, or buffer accessors if that model is used. inside the kernel.

Example (buffer/accessor model):

q.submit([&](handler& cgh) {
      accessor in_accessor(in_buf, cgh, read_only);
      accessor out_accessor(out_buf, cgh, write_only);
        cgh.single_task<NoAliases>([=]() [[intel::kernel_args_restrict]] {
          for (int i = 0; i < N; i++)
              out_accessor[i] = in_accessor[i];
        });
});

On CUDA platforms, data that is constant for at least the lifetime of a kernel can be cached in the texture cache.

This can be achieved by using the sycl::ext::oneapi::experimental::cuda::ldg function. It takes a pointer to device memory and returns the value stored at that address by loading it through the L1/tex cache. For example:

float some_value = ldg(&some_data_in_device_memory[some_index]);

There are many factors impacting the performance benefit of using the texture cache; as such it can be rather difficult to achieve the biggest possible speed-up. Many use-cases in fact receive little to no benefit. However, as performance degradation is both unlikely and small in magnitude when it does happen, and using ldg requires minimal code changes, it can be a great way to quickly improve kernel performance.

Note that ldg can be also be portably used on all other platforms including HIP AMD. However, CUDA is currently the only platform where ldg usage will lead to special caching. The HIP AMD backend always loads all register data into both the L1 and L2 caches whether or not ldg is used.

A more detailed look into the texture cache can be found in this blog post, and the exact detail about the ldg function can be found in the corresponding extension documentation.