Known issues#

HIP#

  • When using HIPManagedSpace, the memory migrates between the CPU and the GPU if:

    • the hardware supports it

    • the kernel was compiled to support page migration

    • the environment variable HSA_XNACK is set to 1

    See here for more explanation.

  • Compatibility issue between HIP and gcc 8. You may encounter the following error:

    error: reference to __host__ function 'operator new' in __host__ __device__ function

    gcc 7, 9, and later do not have this issue.

Mathematical functions#

  • Compatibilty issue with using-directives and mathematical functions:

#include <Kokkos_Core.hpp>

using namespace Kokkos;  // avoid using-directives

KOKKOS_FUNCTION void do_math() {
  auto sqrt5 = sqrt(5);  // error: ambiguous ::sqrt or Kokkos::sqrt?
}

The using-directive using namespace Kokkos; is highly discouraged (see Kokkos compatibility guidelines) and will cause compilation errors in presence of unqualified calls to mathematical functions. Instead, prefer explicit qualification Kokkos::sqrt or an using-declaration using Kokkos::sqrt; at local scope.

Mathematical constants#

  • Avoid taking the address of mathematical constants in device code. It is not supported by some toolchains, hence not portable.

#include <Kokkos_Core.hpp>

KOKKOS_FUNCTION void do_math() {
  // complex constructor takes scalar arguments by reference!
  Kokkos::complex z1(Kokkos::numbers::pi);
  // error: identifier "Kokkos::numbers::pi" is undefined in device code

  // 1*pi is a temporary
  Kokkos::complex z2(1 * Kokkos::numbers::pi);  // OK

  // copy into a local variable
  auto pi = Kokkos::numbers::pi;
  Kokkos::complex z3(pi);  // OK
}