Enable building to target AMD GPUs

CMakeLists.txt

@@ -29,25 +29,64 @@ option(DPCTL_TARGET_CUDA
     "Build DPCTL to target CUDA devices"
     OFF
 )
+set(DPCTL_TARGET_HIP
+    ""
+    CACHE STRING
+    "Build DPCTL to target a HIP device architecture"
+)
 option(DPCTL_WITH_REDIST "Build DPCTL assuming DPC++ redistributable is installed into Python prefix" OFF)
 
 find_package(IntelSYCL REQUIRED PATHS ${CMAKE_SOURCE_DIR}/cmake NO_DEFAULT_PATH)
 
+set(_dpctl_sycl_target_compile_options)
+set(_dpctl_sycl_target_link_options)
+
 set(_dpctl_sycl_targets)
+set(_dpctl_amd_targets)
 if ("x${DPCTL_SYCL_TARGETS}" STREQUAL "x")
-   if(DPCTL_TARGET_CUDA)
+   if (DPCTL_TARGET_CUDA)
       set(_dpctl_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
    else()
-      if(DEFINED ENV{DPCTL_TARGET_CUDA})
+      if (DEFINED ENV{DPCTL_TARGET_CUDA})
           set(_dpctl_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
       endif()
    endif()
+   if (NOT "x${DPCTL_TARGET_HIP}" STREQUAL "x")
+      set(_dpctl_amd_targets ${DPCTL_TARGET_HIP})
+      if(_dpctl_sycl_targets)
+         set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,${_dpctl_sycl_targets}")
+      else()
+         set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,spir64-unknown-unknown")
+      endif()
+   else()
+      if (DEFINED ENV{DPCTL_TARGET_HIP})
+         set(_dpctl_amd_targets $ENV{DPCTL_TARGET_HIP})
+         if(_dpctl_sycl_targets)
+            set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,${_dpctl_sycl_targets}")
+         else()
+            set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,spir64-unknown-unknown")
+         endif()
+      endif()
+   endif()
 else()
    set(_dpctl_sycl_targets ${DPCTL_SYCL_TARGETS})
+   if (NOT "x${DPCTL_TARGET_HIP}" STREQUAL "x")
+      set(_dpctl_amd_targets ${DPCTL_TARGET_HIP})
+   else()
+      if (DEFINED ENV{DPCTL_TARGET_HIP})
+         set(_dpctl_amd_targets $ENV{DPCTL_TARGET_HIP})
+      endif()
+   endif()
 endif()
 
-if(_dpctl_sycl_targets)
+if (_dpctl_sycl_targets)
    message(STATUS "Compiling for -fsycl-targets=${_dpctl_sycl_targets}")
+   list(APPEND _dpctl_sycl_target_compile_options -fsycl-targets=${_dpctl_sycl_targets})
+   list(APPEND _dpctl_sycl_target_link_options -fsycl-targets=${_dpctl_sycl_targets})
+   if(_dpctl_amd_targets)
+      list(APPEND _dpctl_sycl_target_compile_options -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=${_dpctl_amd_targets})
+      list(APPEND _dpctl_sycl_target_link_options -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=${_dpctl_amd_targets})
+   endif()
 endif()
 
 add_subdirectory(libsyclinterface)

docs/doc_sources/beginners_guides/installation.rst

@@ -158,12 +158,28 @@ The following plugins from CodePlay are supported:
 .. _codeplay_nv_plugin: https://developer.codeplay.com/products/oneapi/nvidia/
 .. _codeplay_amd_plugin: https://developer.codeplay.com/products/oneapi/amd/
 
-Build ``dpctl`` as follows:
+``dpctl`` can be built for CUDA devices as follows:
 
 .. code-block:: bash
 
     python scripts/build_locally.py --verbose --cmake-opts="-DDPCTL_TARGET_CUDA=ON"
 
+And for AMD devices
+
+.. code-block:: bash
+
+    python scripts/build_locally.py --verbose --cmake-opts="-DDPCTL_TARGET_HIP=gfx1030"
+
+Note that the  `oneAPI for AMD GPUs` plugin requires the architecture be specified and only
+one architecture can be specified at a time.
+
+It is, however, possible to build for Intel devices, CUDA devices, and an AMD device
+architecture all at once:
+
+.. code-block:: bash
+
+    python scripts/build_locally.py --verbose --cmake-opts="-DDPCTL_TARGET_CUDA=ON \
+    -DDPCTL_TARGET_HIP=gfx1030"
 
 Running Examples and Tests
 ==========================

docs/doc_sources/beginners_guides/managing_devices.rst

@@ -123,6 +123,9 @@ of valid settings are:
     * - ``cuda:*``
       - All devices only from CUDA backend are available
 
+    * - ``hip:*``
+      - All devices only from HIP backend are available
+
     * - ``level_zero:0,1``
       - Two specific devices from Level-Zero backend are available
 

dpctl/CMakeLists.txt

@@ -116,12 +116,12 @@ function(build_dpctl_ext _trgt _src _dest)
             target_compile_options(
                 ${_trgt}
                 PRIVATE
-                -fsycl-targets=${_dpctl_sycl_targets}
+                ${_dpctl_sycl_target_compile_options}
             )
             target_link_options(
                 ${_trgt}
                 PRIVATE
-                -fsycl-targets=${_dpctl_sycl_targets}
+                ${_dpctl_sycl_target_link_options}
             )
         endif()
     endif()

dpctl/_backend.pxd

@@ -43,6 +43,7 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h":
     ctypedef enum _backend_type 'DPCTLSyclBackendType':
         _ALL_BACKENDS    'DPCTL_ALL_BACKENDS'
         _CUDA            'DPCTL_CUDA'
+        _HIP             'DPCTL_HIP'
         _LEVEL_ZERO      'DPCTL_LEVEL_ZERO'
         _OPENCL          'DPCTL_OPENCL'
         _UNKNOWN_BACKEND 'DPCTL_UNKNOWN_BACKEND'

dpctl/_sycl_device.pyx

@@ -172,6 +172,8 @@ cdef list _get_devices(DPCTLDeviceVectorRef DVRef):
 cdef str _backend_type_to_filter_string_part(_backend_type BTy):
     if BTy == _backend_type._CUDA:
         return "cuda"
+    elif BTy == _backend_type._HIP:
+        return "hip"
     elif BTy == _backend_type._LEVEL_ZERO:
         return "level_zero"
     elif BTy == _backend_type._OPENCL:
@@ -425,6 +427,8 @@ cdef class SyclDevice(_SyclDevice):
         )
         if BTy == _backend_type._CUDA:
             return backend_type.cuda
+        elif BTy == _backend_type._HIP:
+            return backend_type.hip
         elif BTy == _backend_type._LEVEL_ZERO:
             return backend_type.level_zero
         elif BTy == _backend_type._OPENCL:

dpctl/_sycl_device_factory.pyx

@@ -71,6 +71,8 @@ cdef _backend_type _string_to_dpctl_sycl_backend_ty(str backend_str):
         return _backend_type._ALL_BACKENDS
     elif backend_str == "cuda":
         return _backend_type._CUDA
+    elif backend_str == "hip":
+        return _backend_type._HIP
     elif backend_str == "level_zero":
         return _backend_type._LEVEL_ZERO
     elif backend_str == "opencl":
@@ -100,6 +102,8 @@ cdef _device_type _string_to_dpctl_sycl_device_ty(str dty_str):
 cdef _backend_type _enum_to_dpctl_sycl_backend_ty(BTy):
     if BTy == backend_type.all:
         return _backend_type._ALL_BACKENDS
+    elif BTy == backend_type.hip:
+        return _backend_type._HIP
     elif BTy == backend_type.cuda:
         return _backend_type._CUDA
     elif BTy == backend_type.level_zero:
@@ -153,7 +157,7 @@ cpdef list get_devices(backend=backend_type.all, device_type=device_type_t.all):
         backend (optional):
             A :class:`dpctl.backend_type` enum value or a string that
             specifies a SYCL backend. Currently, accepted values are: "cuda",
-            "opencl", "level_zero", or "all".
+            "hip", "opencl", "level_zero", or "all".
             Default: ``dpctl.backend_type.all``.
         device_type (optional):
             A :class:`dpctl.device_type` enum value or a string that
@@ -209,7 +213,7 @@ cpdef int get_num_devices(
         backend (optional):
             A :class:`dpctl.backend_type` enum value or a string that
             specifies a SYCL backend. Currently, accepted values are: "cuda",
-            "opencl", "level_zero", or "all".
+            "hip", "opencl", "level_zero", or "all".
             Default: ``dpctl.backend_type.all``.
         device_type (optional):
             A :class:`dpctl.device_type` enum value or a string that

dpctl/_sycl_event.pyx

@@ -333,6 +333,8 @@ cdef class SyclEvent(_SyclEvent):
             return backend_type.level_zero
         elif BE == _backend_type._CUDA:
             return backend_type.cuda
+        elif BE == _backend_type._HIP:
+            return backend_type.hip
         else:
             raise ValueError("Unknown backend type.")
 

dpctl/_sycl_platform.pyx

@@ -294,6 +294,8 @@ cdef class SyclPlatform(_SyclPlatform):
         )
         if BTy == _backend_type._CUDA:
             return backend_type.cuda
+        elif BTy == _backend_type._HIP:
+            return backend_type.hip
         elif BTy == _backend_type._LEVEL_ZERO:
             return backend_type.level_zero
         elif BTy == _backend_type._OPENCL:

dpctl/_sycl_queue.pyx

@@ -886,6 +886,8 @@ cdef class SyclQueue(_SyclQueue):
             return backend_type.level_zero
         elif BE == _backend_type._CUDA:
             return backend_type.cuda
+        elif BE == _backend_type._HIP:
+            return backend_type.hip
         else:
             raise ValueError("Unknown backend type.")
 

dpctl/enum_types.py

@@ -69,6 +69,7 @@ class backend_type(Enum):
 
     all = auto()
     cuda = auto()
+    hip = auto()
     level_zero = auto()
     opencl = auto()
 

dpctl/tensor/CMakeLists.txt

@@ -300,12 +300,12 @@ foreach(python_module_name ${_py_trgts})
         target_compile_options(
             ${python_module_name}
             PRIVATE
-            -fsycl-targets=${_dpctl_sycl_targets}
+            ${_dpctl_sycl_target_compile_options}
         )
         target_link_options(
             ${python_module_name}
             PRIVATE
-            -fsycl-targets=${_dpctl_sycl_targets}
+            ${_dpctl_sycl_target_link_options}
         )
     endif()
     # TODO: update source so they refernece individual libraries instead of

dpctl/tensor/libtensor/include/utils/math_utils.hpp

@@ -126,30 +126,10 @@ template <typename T> T logaddexp(T x, T y)
         const T tmp = x - y;
         constexpr T zero(0);
 
-        if constexpr (std::is_same_v<T, sycl::half>) {
-            return (tmp > zero)
-                       ? (x + sycl::log1p(sycl::exp(-tmp)))
-                       : ((tmp <= zero) ? y + sycl::log1p(sycl::exp(tmp))
-                                        : std::numeric_limits<T>::quiet_NaN());
-        }
-        else {
-            if constexpr (std::is_same_v<T, double>) {
-                // FIXME: switch to `sycl::log1p` when
-                // compiler segfault in CUDA build is fixed
-                return (tmp > zero)
-                           ? (x + std::log1p(sycl::exp(-tmp)))
-                           : ((tmp <= zero)
-                                  ? y + std::log1p(sycl::exp(tmp))
-                                  : std::numeric_limits<T>::quiet_NaN());
-            }
-            else {
-                return (tmp > zero)
-                           ? (x + sycl::log1p(sycl::exp(-tmp)))
-                           : ((tmp <= zero)
-                                  ? y + sycl::log1p(sycl::exp(tmp))
-                                  : std::numeric_limits<T>::quiet_NaN());
-            }
-        }
+        return (tmp > zero)
+                   ? (x + sycl::log1p(sycl::exp(-tmp)))
+                   : ((tmp <= zero) ? y + sycl::log1p(sycl::exp(tmp))
+                                    : std::numeric_limits<T>::quiet_NaN());
     }
 }
 

dpctl/tests/test_sycl_device_factory.py

@@ -60,6 +60,8 @@ def string_to_device_type(dty_str):
 def string_to_backend_type(bty_str):
     if bty_str == "cuda":
         return bty.cuda
+    elif bty_str == "hip":
+        return bty.hip
     elif bty_str == "host":
         return bty.host
     elif bty_str == "level_zero":

dpctl/utils/CMakeLists.txt

@@ -50,12 +50,12 @@ foreach(python_module_name ${_pybind11_targets})
         target_compile_options(
             ${python_module_name}
             PRIVATE
-            -fsycl-targets=${_dpctl_sycl_targets}
+            ${_dpctl_sycl_target_compile_options}
         )
         target_link_options(
             ${python_module_name}
             PRIVATE
-            -fsycl-targets=${_dpctl_sycl_targets}
+            ${_dpctl_sycl_target_link_options}
         )
     endif()
     # TODO: update source so they refernece individual libraries instead of

libsyclinterface/CMakeLists.txt

@@ -244,13 +244,13 @@ if(_dpctl_sycl_targets)
     target_compile_options(
         DPCTLSyclInterface
         PRIVATE
-        -fsycl-targets=${_dpctl_sycl_targets}
+        ${_dpctl_sycl_target_compile_options}
     )
     target_link_options(
         DPCTLSyclInterface
         PRIVATE
-        -fsycl-targets=${_dpctl_sycl_targets}
-   )
+        ${_dpctl_sycl_target_link_options}
+    )
 endif()
 
 if(DPCTL_GENERATE_COVERAGE)

libsyclinterface/helper/source/dpctl_utils_helper.cpp

@@ -94,6 +94,8 @@ backend DPCTL_DPCTLBackendTypeToSyclBackend(DPCTLSyclBackendType BeTy)
         return backend::opencl;
     case DPCTLSyclBackendType::DPCTL_ALL_BACKENDS:
         return backend::all;
+    case DPCTLSyclBackendType::DPCTL_HIP:
+        return backend::ext_oneapi_hip;
     default:
         throw std::runtime_error("Unsupported backend type");
     }
@@ -108,6 +110,8 @@ DPCTLSyclBackendType DPCTL_SyclBackendToDPCTLBackendType(backend B)
         return DPCTLSyclBackendType::DPCTL_LEVEL_ZERO;
     case backend::opencl:
         return DPCTLSyclBackendType::DPCTL_OPENCL;
+    case backend::ext_oneapi_hip:
+        return DPCTLSyclBackendType::DPCTL_HIP;
     default:
         return DPCTLSyclBackendType::DPCTL_UNKNOWN_BACKEND;
     }
@@ -467,6 +471,9 @@ std::string DPCTL_GetDeviceFilterString(const device &Device)
     case backend::opencl:
         ss << "opencl";
         break;
+    case backend::ext_oneapi_hip:
+        ss << "hip";
+        break;
     default:
         ss << "unknown";
     };

libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h

@@ -54,7 +54,8 @@ typedef enum
     DPCTL_LEVEL_ZERO      = 1 << 17,
     DPCTL_OPENCL          = 1 << 18,
     DPCTL_UNKNOWN_BACKEND = 0,
-    DPCTL_ALL_BACKENDS    = ((1<<5)-1) << 16
+    DPCTL_ALL_BACKENDS    = ((1<<5)-1) << 16,
+    DPCTL_HIP             = 1 << 19,
     // clang-format on
 } DPCTLSyclBackendType;
 

libsyclinterface/source/dpctl_sycl_context_interface.cpp

@@ -192,6 +192,8 @@ DPCTLContext_GetBackend(__dpctl_keep const DPCTLSyclContextRef CtxRef)
         return DPCTL_LEVEL_ZERO;
     case backend::ext_oneapi_cuda:
         return DPCTL_CUDA;
+    case backend::ext_oneapi_hip:
+        return DPCTL_HIP;
     default:
         return DPCTL_UNKNOWN_BACKEND;
     }

libsyclinterface/tests/CMakeLists.txt

@@ -55,16 +55,15 @@ add_sycl_to_target(
 )
 
 if(_dpctl_sycl_targets)
-    # make fat binary
     target_compile_options(
         dpctl_c_api_tests
         PRIVATE
-        -fsycl-targets=${_dpctl_sycl_targets}
+        ${_dpctl_sycl_target_compile_options}
     )
     target_link_options(
         dpctl_c_api_tests
         PRIVATE
-        -fsycl-targets=${_dpctl_sycl_targets}
+        ${_dpctl_sycl_target_link_options}
     )
 endif()
 
@@ -85,16 +84,15 @@ target_include_directories(dpctl_c_api_tests
 )
 
 if(_dpctl_sycl_targets)
-    # make fat binary
     target_compile_options(
         dpctl_c_api_tests
         PRIVATE
-        -fsycl-targets=${_dpctl_sycl_targets}
+        ${_dpctl_sycl_target_compile_options}
     )
     target_link_options(
         dpctl_c_api_tests
         PRIVATE
-        -fsycl-targets=${_dpctl_sycl_targets}
+        ${_dpctl_sycl_target_link_options}
     )
 endif()