Skip to content

Commit eefc82b

Browse files
ndgrigorianndgrigorian
authored
Merge pull request #1731 from IntelPython/feature/enable-amd-builds
Enable building to target AMD GPUs
2 parents ea6ae0b + c316380 commit eefc82b

26 files changed

+128
-46
lines changed

‎CMakeLists.txt

Lines changed: 42 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -29,25 +29,64 @@ option(DPCTL_TARGET_CUDA
2929
"Build DPCTL to target CUDA devices"
3030
OFF
3131
)
32+
set(DPCTL_TARGET_HIP
33+
""
34+
CACHE STRING
35+
"Build DPCTL to target a HIP device architecture"
36+
)
3237
option(DPCTL_WITH_REDIST "Build DPCTL assuming DPC++ redistributable is installed into Python prefix" OFF)
3338

3439
find_package(IntelSYCL REQUIRED PATHS ${CMAKE_SOURCE_DIR}/cmake NO_DEFAULT_PATH)
3540

41+
set(_dpctl_sycl_target_compile_options)
42+
set(_dpctl_sycl_target_link_options)
43+
3644
set(_dpctl_sycl_targets)
45+
set(_dpctl_amd_targets)
3746
if ("x${DPCTL_SYCL_TARGETS}" STREQUAL "x")
38-
if(DPCTL_TARGET_CUDA)
47+
if (DPCTL_TARGET_CUDA)
3948
set(_dpctl_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
4049
else()
41-
if(DEFINED ENV{DPCTL_TARGET_CUDA})
50+
if (DEFINED ENV{DPCTL_TARGET_CUDA})
4251
set(_dpctl_sycl_targets "nvptx64-nvidia-cuda,spir64-unknown-unknown")
4352
endif()
4453
endif()
54+
if (NOT "x${DPCTL_TARGET_HIP}" STREQUAL "x")
55+
set(_dpctl_amd_targets ${DPCTL_TARGET_HIP})
56+
if(_dpctl_sycl_targets)
57+
set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,${_dpctl_sycl_targets}")
58+
else()
59+
set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,spir64-unknown-unknown")
60+
endif()
61+
else()
62+
if (DEFINED ENV{DPCTL_TARGET_HIP})
63+
set(_dpctl_amd_targets $ENV{DPCTL_TARGET_HIP})
64+
if(_dpctl_sycl_targets)
65+
set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,${_dpctl_sycl_targets}")
66+
else()
67+
set(_dpctl_sycl_targets "amdgcn-amd-amdhsa,spir64-unknown-unknown")
68+
endif()
69+
endif()
70+
endif()
4571
else()
4672
set(_dpctl_sycl_targets ${DPCTL_SYCL_TARGETS})
73+
if (NOT "x${DPCTL_TARGET_HIP}" STREQUAL "x")
74+
set(_dpctl_amd_targets ${DPCTL_TARGET_HIP})
75+
else()
76+
if (DEFINED ENV{DPCTL_TARGET_HIP})
77+
set(_dpctl_amd_targets $ENV{DPCTL_TARGET_HIP})
78+
endif()
79+
endif()
4780
endif()
4881

49-
if(_dpctl_sycl_targets)
82+
if (_dpctl_sycl_targets)
5083
message(STATUS "Compiling for -fsycl-targets=${_dpctl_sycl_targets}")
84+
list(APPEND _dpctl_sycl_target_compile_options -fsycl-targets=${_dpctl_sycl_targets})
85+
list(APPEND _dpctl_sycl_target_link_options -fsycl-targets=${_dpctl_sycl_targets})
86+
if(_dpctl_amd_targets)
87+
list(APPEND _dpctl_sycl_target_compile_options -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=${_dpctl_amd_targets})
88+
list(APPEND _dpctl_sycl_target_link_options -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=${_dpctl_amd_targets})
89+
endif()
5190
endif()
5291

5392
add_subdirectory(libsyclinterface)

‎docs/doc_sources/beginners_guides/installation.rst

Lines changed: 17 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -158,12 +158,28 @@ The following plugins from CodePlay are supported:
158158
.. _codeplay_nv_plugin: https://developer.codeplay.com/products/oneapi/nvidia/
159159
.. _codeplay_amd_plugin: https://developer.codeplay.com/products/oneapi/amd/
160160

161-
Build ``dpctl`` as follows:
161+
``dpctl`` can be built for CUDA devices as follows:
162162

163163
.. code-block:: bash
164164
165165
python scripts/build_locally.py --verbose --cmake-opts="-DDPCTL_TARGET_CUDA=ON"
166166
167+
And for AMD devices
168+
169+
.. code-block:: bash
170+
171+
python scripts/build_locally.py --verbose --cmake-opts="-DDPCTL_TARGET_HIP=gfx1030"
172+
173+
Note that the `oneAPI for AMD GPUs` plugin requires the architecture be specified and only
174+
one architecture can be specified at a time.
175+
176+
It is, however, possible to build for Intel devices, CUDA devices, and an AMD device
177+
architecture all at once:
178+
179+
.. code-block:: bash
180+
181+
python scripts/build_locally.py --verbose --cmake-opts="-DDPCTL_TARGET_CUDA=ON \
182+
-DDPCTL_TARGET_HIP=gfx1030"
167183
168184
Running Examples and Tests
169185
==========================

‎docs/doc_sources/beginners_guides/managing_devices.rst

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -123,6 +123,9 @@ of valid settings are:
123123
* - ``cuda:*``
124124
- All devices only from CUDA backend are available
125125
126+
* - ``hip:*``
127+
- All devices only from HIP backend are available
128+
126129
* - ``level_zero:0,1``
127130
- Two specific devices from Level-Zero backend are available
128131

‎dpctl/CMakeLists.txt

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -116,12 +116,12 @@ function(build_dpctl_ext _trgt _src _dest)
116116
target_compile_options(
117117
${_trgt}
118118
PRIVATE
119-
-fsycl-targets=${_dpctl_sycl_targets}
119+
${_dpctl_sycl_target_compile_options}
120120
)
121121
target_link_options(
122122
${_trgt}
123123
PRIVATE
124-
-fsycl-targets=${_dpctl_sycl_targets}
124+
${_dpctl_sycl_target_link_options}
125125
)
126126
endif()
127127
endif()

‎dpctl/_backend.pxd

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,7 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h":
4343
ctypedef enum _backend_type 'DPCTLSyclBackendType':
4444
_ALL_BACKENDS 'DPCTL_ALL_BACKENDS'
4545
_CUDA 'DPCTL_CUDA'
46+
_HIP 'DPCTL_HIP'
4647
_LEVEL_ZERO 'DPCTL_LEVEL_ZERO'
4748
_OPENCL 'DPCTL_OPENCL'
4849
_UNKNOWN_BACKEND 'DPCTL_UNKNOWN_BACKEND'

‎dpctl/_sycl_device.pyx

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -172,6 +172,8 @@ cdef list _get_devices(DPCTLDeviceVectorRef DVRef):
172172
cdef str _backend_type_to_filter_string_part(_backend_type BTy):
173173
if BTy == _backend_type._CUDA:
174174
return "cuda"
175+
elif BTy == _backend_type._HIP:
176+
return "hip"
175177
elif BTy == _backend_type._LEVEL_ZERO:
176178
return "level_zero"
177179
elif BTy == _backend_type._OPENCL:
@@ -425,6 +427,8 @@ cdef class SyclDevice(_SyclDevice):
425427
)
426428
if BTy == _backend_type._CUDA:
427429
return backend_type.cuda
430+
elif BTy == _backend_type._HIP:
431+
return backend_type.hip
428432
elif BTy == _backend_type._LEVEL_ZERO:
429433
return backend_type.level_zero
430434
elif BTy == _backend_type._OPENCL:

‎dpctl/_sycl_device_factory.pyx

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -71,6 +71,8 @@ cdef _backend_type _string_to_dpctl_sycl_backend_ty(str backend_str):
7171
return _backend_type._ALL_BACKENDS
7272
elif backend_str == "cuda":
7373
return _backend_type._CUDA
74+
elif backend_str == "hip":
75+
return _backend_type._HIP
7476
elif backend_str == "level_zero":
7577
return _backend_type._LEVEL_ZERO
7678
elif backend_str == "opencl":
@@ -100,6 +102,8 @@ cdef _device_type _string_to_dpctl_sycl_device_ty(str dty_str):
100102
cdef _backend_type _enum_to_dpctl_sycl_backend_ty(BTy):
101103
if BTy == backend_type.all:
102104
return _backend_type._ALL_BACKENDS
105+
elif BTy == backend_type.hip:
106+
return _backend_type._HIP
103107
elif BTy == backend_type.cuda:
104108
return _backend_type._CUDA
105109
elif BTy == backend_type.level_zero:
@@ -153,7 +157,7 @@ cpdef list get_devices(backend=backend_type.all, device_type=device_type_t.all):
153157
backend (optional):
154158
A :class:`dpctl.backend_type` enum value or a string that
155159
specifies a SYCL backend. Currently, accepted values are: "cuda",
156-
"opencl", "level_zero", or "all".
160+
"hip", "opencl", "level_zero", or "all".
157161
Default: ``dpctl.backend_type.all``.
158162
device_type (optional):
159163
A :class:`dpctl.device_type` enum value or a string that
@@ -209,7 +213,7 @@ cpdef int get_num_devices(
209213
backend (optional):
210214
A :class:`dpctl.backend_type` enum value or a string that
211215
specifies a SYCL backend. Currently, accepted values are: "cuda",
212-
"opencl", "level_zero", or "all".
216+
"hip", "opencl", "level_zero", or "all".
213217
Default: ``dpctl.backend_type.all``.
214218
device_type (optional):
215219
A :class:`dpctl.device_type` enum value or a string that

‎dpctl/_sycl_event.pyx

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -333,6 +333,8 @@ cdef class SyclEvent(_SyclEvent):
333333
return backend_type.level_zero
334334
elif BE == _backend_type._CUDA:
335335
return backend_type.cuda
336+
elif BE == _backend_type._HIP:
337+
return backend_type.hip
336338
else:
337339
raise ValueError("Unknown backend type.")
338340

‎dpctl/_sycl_platform.pyx

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -294,6 +294,8 @@ cdef class SyclPlatform(_SyclPlatform):
294294
)
295295
if BTy == _backend_type._CUDA:
296296
return backend_type.cuda
297+
elif BTy == _backend_type._HIP:
298+
return backend_type.hip
297299
elif BTy == _backend_type._LEVEL_ZERO:
298300
return backend_type.level_zero
299301
elif BTy == _backend_type._OPENCL:

‎dpctl/_sycl_queue.pyx

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -886,6 +886,8 @@ cdef class SyclQueue(_SyclQueue):
886886
return backend_type.level_zero
887887
elif BE == _backend_type._CUDA:
888888
return backend_type.cuda
889+
elif BE == _backend_type._HIP:
890+
return backend_type.hip
889891
else:
890892
raise ValueError("Unknown backend type.")
891893

‎dpctl/enum_types.py

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -69,6 +69,7 @@ class backend_type(Enum):
6969

7070
all = auto()
7171
cuda = auto()
72+
hip = auto()
7273
level_zero = auto()
7374
opencl = auto()
7475

‎dpctl/tensor/CMakeLists.txt

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -300,12 +300,12 @@ foreach(python_module_name ${_py_trgts})
300300
target_compile_options(
301301
${python_module_name}
302302
PRIVATE
303-
-fsycl-targets=${_dpctl_sycl_targets}
303+
${_dpctl_sycl_target_compile_options}
304304
)
305305
target_link_options(
306306
${python_module_name}
307307
PRIVATE
308-
-fsycl-targets=${_dpctl_sycl_targets}
308+
${_dpctl_sycl_target_link_options}
309309
)
310310
endif()
311311
# TODO: update source so they refernece individual libraries instead of

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

Lines changed: 4 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -126,30 +126,10 @@ template <typename T> T logaddexp(T x, T y)
126126
const T tmp = x - y;
127127
constexpr T zero(0);
128128

129-
if constexpr (std::is_same_v<T, sycl::half>) {
130-
return (tmp > zero)
131-
? (x + sycl::log1p(sycl::exp(-tmp)))
132-
: ((tmp <= zero) ? y + sycl::log1p(sycl::exp(tmp))
133-
: std::numeric_limits<T>::quiet_NaN());
134-
}
135-
else {
136-
if constexpr (std::is_same_v<T, double>) {
137-
// FIXME: switch to `sycl::log1p` when
138-
// compiler segfault in CUDA build is fixed
139-
return (tmp > zero)
140-
? (x + std::log1p(sycl::exp(-tmp)))
141-
: ((tmp <= zero)
142-
? y + std::log1p(sycl::exp(tmp))
143-
: std::numeric_limits<T>::quiet_NaN());
144-
}
145-
else {
146-
return (tmp > zero)
147-
? (x + sycl::log1p(sycl::exp(-tmp)))
148-
: ((tmp <= zero)
149-
? y + sycl::log1p(sycl::exp(tmp))
150-
: std::numeric_limits<T>::quiet_NaN());
151-
}
152-
}
129+
return (tmp > zero)
130+
? (x + sycl::log1p(sycl::exp(-tmp)))
131+
: ((tmp <= zero) ? y + sycl::log1p(sycl::exp(tmp))
132+
: std::numeric_limits<T>::quiet_NaN());
153133
}
154134
}
155135

‎dpctl/tests/test_sycl_device_factory.py

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -60,6 +60,8 @@ def string_to_device_type(dty_str):
6060
def string_to_backend_type(bty_str):
6161
if bty_str == "cuda":
6262
return bty.cuda
63+
elif bty_str == "hip":
64+
return bty.hip
6365
elif bty_str == "host":
6466
return bty.host
6567
elif bty_str == "level_zero":

‎dpctl/utils/CMakeLists.txt

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -50,12 +50,12 @@ foreach(python_module_name ${_pybind11_targets})
5050
target_compile_options(
5151
${python_module_name}
5252
PRIVATE
53-
-fsycl-targets=${_dpctl_sycl_targets}
53+
${_dpctl_sycl_target_compile_options}
5454
)
5555
target_link_options(
5656
${python_module_name}
5757
PRIVATE
58-
-fsycl-targets=${_dpctl_sycl_targets}
58+
${_dpctl_sycl_target_link_options}
5959
)
6060
endif()
6161
# TODO: update source so they refernece individual libraries instead of

‎libsyclinterface/CMakeLists.txt

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -244,13 +244,13 @@ if(_dpctl_sycl_targets)
244244
target_compile_options(
245245
DPCTLSyclInterface
246246
PRIVATE
247-
-fsycl-targets=${_dpctl_sycl_targets}
247+
${_dpctl_sycl_target_compile_options}
248248
)
249249
target_link_options(
250250
DPCTLSyclInterface
251251
PRIVATE
252-
-fsycl-targets=${_dpctl_sycl_targets}
253-
)
252+
${_dpctl_sycl_target_link_options}
253+
)
254254
endif()
255255

256256
if(DPCTL_GENERATE_COVERAGE)

‎libsyclinterface/helper/source/dpctl_utils_helper.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -94,6 +94,8 @@ backend DPCTL_DPCTLBackendTypeToSyclBackend(DPCTLSyclBackendType BeTy)
9494
return backend::opencl;
9595
case DPCTLSyclBackendType::DPCTL_ALL_BACKENDS:
9696
return backend::all;
97+
case DPCTLSyclBackendType::DPCTL_HIP:
98+
return backend::ext_oneapi_hip;
9799
default:
98100
throw std::runtime_error("Unsupported backend type");
99101
}
@@ -108,6 +110,8 @@ DPCTLSyclBackendType DPCTL_SyclBackendToDPCTLBackendType(backend B)
108110
return DPCTLSyclBackendType::DPCTL_LEVEL_ZERO;
109111
case backend::opencl:
110112
return DPCTLSyclBackendType::DPCTL_OPENCL;
113+
case backend::ext_oneapi_hip:
114+
return DPCTLSyclBackendType::DPCTL_HIP;
111115
default:
112116
return DPCTLSyclBackendType::DPCTL_UNKNOWN_BACKEND;
113117
}
@@ -467,6 +471,9 @@ std::string DPCTL_GetDeviceFilterString(const device &Device)
467471
case backend::opencl:
468472
ss << "opencl";
469473
break;
474+
case backend::ext_oneapi_hip:
475+
ss << "hip";
476+
break;
470477
default:
471478
ss << "unknown";
472479
};

‎libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,8 @@ typedef enum
5454
DPCTL_LEVEL_ZERO = 1 << 17,
5555
DPCTL_OPENCL = 1 << 18,
5656
DPCTL_UNKNOWN_BACKEND = 0,
57-
DPCTL_ALL_BACKENDS = ((1<<5)-1) << 16
57+
DPCTL_ALL_BACKENDS = ((1<<5)-1) << 16,
58+
DPCTL_HIP = 1 << 19,
5859
// clang-format on
5960
} DPCTLSyclBackendType;
6061

‎libsyclinterface/source/dpctl_sycl_context_interface.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -192,6 +192,8 @@ DPCTLContext_GetBackend(__dpctl_keep const DPCTLSyclContextRef CtxRef)
192192
return DPCTL_LEVEL_ZERO;
193193
case backend::ext_oneapi_cuda:
194194
return DPCTL_CUDA;
195+
case backend::ext_oneapi_hip:
196+
return DPCTL_HIP;
195197
default:
196198
return DPCTL_UNKNOWN_BACKEND;
197199
}

‎libsyclinterface/tests/CMakeLists.txt

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -55,16 +55,15 @@ add_sycl_to_target(
5555
)
5656

5757
if(_dpctl_sycl_targets)
58-
# make fat binary
5958
target_compile_options(
6059
dpctl_c_api_tests
6160
PRIVATE
62-
-fsycl-targets=${_dpctl_sycl_targets}
61+
${_dpctl_sycl_target_compile_options}
6362
)
6463
target_link_options(
6564
dpctl_c_api_tests
6665
PRIVATE
67-
-fsycl-targets=${_dpctl_sycl_targets}
66+
${_dpctl_sycl_target_link_options}
6867
)
6968
endif()
7069

@@ -85,16 +84,15 @@ target_include_directories(dpctl_c_api_tests
8584
)
8685

8786
if(_dpctl_sycl_targets)
88-
# make fat binary
8987
target_compile_options(
9088
dpctl_c_api_tests
9189
PRIVATE
92-
-fsycl-targets=${_dpctl_sycl_targets}
90+
${_dpctl_sycl_target_compile_options}
9391
)
9492
target_link_options(
9593
dpctl_c_api_tests
9694
PRIVATE
97-
-fsycl-targets=${_dpctl_sycl_targets}
95+
${_dpctl_sycl_target_link_options}
9896
)
9997
endif()
10098

0 commit comments

Comments
 (0)