Skip to content

Development milestone 0.14.6dev4 #1354

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 44 commits into from
Aug 18, 2023
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
44 commits
Select commit Hold shift + click to select a range
7e9f857
Transition sum-reduction from nd_range<2> to nd_range<1>
oleksandr-pavlyk Jul 25, 2023
2f79acb
Replaced nd_range<2> with nd_range<1> for non-atomic case as well
oleksandr-pavlyk Jul 28, 2023
5c4f980
Add test based on example from @ndgrigorian's feedback to PR
oleksandr-pavlyk Jul 28, 2023
a5aee5b
Boolean reductions transitioned from nd_range<2> to nd_range<1>
ndgrigorian Jul 30, 2023
da1cc5b
Strided boolean reduction loop tweak
ndgrigorian Jul 30, 2023
63b2799
Merge remote-tracking branch 'origin/master' into reduction-changes
oleksandr-pavlyk Aug 11, 2023
9f54428
Specify name for the atomic reduction initialization kernel
oleksandr-pavlyk Aug 13, 2023
48c2ad2
Clean up of operator special methods
oleksandr-pavlyk Aug 14, 2023
7d9974e
Removed leftover include iostream
oleksandr-pavlyk Aug 15, 2023
5f298e6
Closes gh-1279
oleksandr-pavlyk Aug 13, 2023
4a2578f
Closes gh-1279 for dpt.sqrt
oleksandr-pavlyk Aug 15, 2023
26862b4
Implements dpctl.tensor.allclose
oleksandr-pavlyk Aug 15, 2023
b121d67
Adds tests for special FP values for dpt.abs and dpt.sqrt
oleksandr-pavlyk Aug 15, 2023
d60d58e
Added tests for type promotion in tensor.allclose
oleksandr-pavlyk Aug 15, 2023
ac54f9e
Use latest sycl bundle to build DPCTL
oleksandr-pavlyk Aug 15, 2023
2dda06c
Run on Ubuntu 22.04, since nightly sycl bundle requires newer GLIBC
oleksandr-pavlyk Aug 15, 2023
9056477
Copy OpenCL loader from oclcpu to compiler's lib
oleksandr-pavlyk Aug 15, 2023
9afdb58
Merge pull request #1346 from IntelPython/use-latest-intel-llvm-sycl-…
oleksandr-pavlyk Aug 15, 2023
0ed60a9
Test environment requires compilers
oleksandr-pavlyk Aug 16, 2023
ee71562
Corrected typo in exception text
antonwolfy Aug 16, 2023
2256fb3
Adding test to prevent gh-1344 from reoccuring
oleksandr-pavlyk Aug 16, 2023
7f7897c
Updated copyright year range
oleksandr-pavlyk Aug 16, 2023
02fd974
Merge pull request #1349 from IntelPython/typo_in_exception
oleksandr-pavlyk Aug 16, 2023
ea6dd27
Fixes per PR review feedback
oleksandr-pavlyk Aug 16, 2023
f36af57
Adds tests for atol/rtol
oleksandr-pavlyk Aug 16, 2023
a75fff8
tensor.allclose to use abs(a-b) < max(atol, rtol*max(abs(a), abs(b)))
oleksandr-pavlyk Aug 16, 2023
47c82f5
Merge pull request #1348 from IntelPython/test-requires-compilers
oleksandr-pavlyk Aug 16, 2023
ff3c680
Completion of fix for gh-1058
oleksandr-pavlyk Aug 16, 2023
53e850d
Closes gh-1350
oleksandr-pavlyk Aug 16, 2023
43b9c06
Added test based on gh-1350
oleksandr-pavlyk Aug 16, 2023
c4312cb
Scale down arguments and scale back the result
oleksandr-pavlyk Aug 16, 2023
9c18949
Merge pull request #1351 from IntelPython/fix-gh-1350
oleksandr-pavlyk Aug 17, 2023
1c09c66
Update README for wheel installation
DenisScherbakov Aug 17, 2023
ec60e9f
Update README.md
oleksandr-pavlyk Aug 17, 2023
ea83961
Update README.md
oleksandr-pavlyk Aug 17, 2023
1468978
Update README.md
oleksandr-pavlyk Aug 17, 2023
bb52bb1
Avoid using sycl::ilogb, but use own implementation
oleksandr-pavlyk Aug 17, 2023
ba9a595
Set defines to use std::abs and std::sqrt on Linux
oleksandr-pavlyk Aug 17, 2023
142190f
Removed stray include iostream
oleksandr-pavlyk Aug 17, 2023
bd996b5
Merge pull request #1353 from IntelPython/fix/install
oleksandr-pavlyk Aug 17, 2023
2f3be1f
Merge pull request #1343 from IntelPython/fix-gh-1279
oleksandr-pavlyk Aug 18, 2023
b008b8b
Merge pull request #1303 from IntelPython/reduction-changes
oleksandr-pavlyk Aug 18, 2023
852f4b1
``dpctl.tensor.where`` output preserves memory order of inputs (#1342)
ndgrigorian Aug 18, 2023
96293fd
Fully enable ``usm_ndarray`` in-place arithmetic operators (#1352)
ndgrigorian Aug 18, 2023
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
25 changes: 13 additions & 12 deletions .github/workflows/os-llvm-sycl-build.yml
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@ on:
jobs:
install-compiler:
name: Build with nightly build of DPC++ toolchain
runs-on: ubuntu-20.04
runs-on: ubuntu-22.04

env:
DOWNLOAD_URL_PREFIX: https://github.com/intel/llvm/releases/download
Expand Down Expand Up @@ -38,16 +38,17 @@ jobs:

- name: Download and install nightly and components
env:
USE_LATEST_SYCLOS: 0
ARTIFACT_NAME: sycl_linux
USE_LATEST_SYCLOS: 1
shell: bash -l {0}
run: |
cd /home/runner/work
mkdir -p sycl_bundle
cd sycl_bundle
if [[ "${USE_LATEST_SYCLOS:-0}" -eq "1" ]]; then
# get list of shas and tags from remote, filter sycl-nightly tags and reverse order
# get list of shas and tags from remote, filter nightly tags and reverse order
export LLVM_TAGS=$(git -c 'versionsort.suffix=-' ls-remote --tags --sort='v:refname' https://github.com/intel/llvm.git | \
grep sycl-nightly | awk '{a[i++]=$0} END {for (j=i-1; j>=0;) print a[j--] }')
grep 'refs/tags/nightly-' | awk '{a[i++]=$0} END {for (j=i-1; j>=0;) print a[j--] }')
# initialize
unset DEPLOY_NIGHTLY_TAG
unset DEPLOY_NIGHTLY_TAG_SHA
Expand All @@ -57,7 +58,7 @@ jobs:
export NEXT_LLVM_TAG_SHA=$(echo ${NEXT_LLVM_TAG} | awk '{print $1}')
export NEXT_NIGHTLY_TAG=$(python3 -c "import sys, urllib.parse as ul; print (ul.quote_plus(sys.argv[1]))" \
$(echo ${NEXT_LLVM_TAG} | awk '{gsub(/^refs\/tags\//, "", $2)} {print $2}'))
if [[ `wget -S --spider ${DOWNLOAD_URL_PREFIX}/${NEXT_NIGHTLY_TAG}/dpcpp-compiler.tar.gz 2>&1 | grep 'HTTP/1.1 200 OK'` ]];
if [[ `wget -S --spider ${DOWNLOAD_URL_PREFIX}/${NEXT_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz 2>&1 | grep 'HTTP/1.1 200 OK'` ]];
then
export DEPLOY_NIGHTLY_TAG=${NEXT_NIGHTLY_TAG}
export DEPLOY_LLVM_TAG_SHA=${NEXT_LLVM_TAG_SHA}
Expand All @@ -77,21 +78,20 @@ jobs:
if [[ -f bundle_id.txt && ( "$(cat bundle_id.txt)" == "${DEPLOY_LLVM_TAG_SHA}" ) ]]; then
echo "Using cached download of ${DEPLOY_LLVM_TAG_SHA}"
else
rm -rf dpcpp-compiler.tar.gz
wget ${DOWNLOAD_URL_PREFIX}/${DEPLOY_NIGHTLY_TAG}/dpcpp-compiler.tar.gz && echo ${DEPLOY_LLVM_TAG_SHA} > bundle_id.txt || rm -rf bundle_id.txt
rm -rf ${ARTIFACT_NAME}.tar.gz
wget ${DOWNLOAD_URL_PREFIX}/${DEPLOY_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz && echo ${DEPLOY_LLVM_TAG_SHA} > bundle_id.txt || rm -rf bundle_id.txt
[ -f ${OCLCPUEXP_FN} ] || wget ${DOWNLOAD_URL_PREFIX}/${DRIVER_PATH}/${OCLCPUEXP_FN} || rm -rf bundle_id.txt
[ -f ${FPGAEMU_FN} ] || wget ${DOWNLOAD_URL_PREFIX}/${DRIVER_PATH}/${FPGAEMU_FN} || rm -rf bundle_id.txt
[ -f ${TBB_FN} ] || wget ${TBB_URL}/${TBB_FN} || rm -rf bundle_id.txt
rm -rf dpcpp_compiler
tar xf dpcpp-compiler.tar.gz
mkdir -p dpcpp_compiler
tar xf ${ARTIFACT_NAME}.tar.gz -C dpcpp_compiler
mkdir -p oclcpuexp
mkdir -p fpgaemu
[ -d oclcpuexp/x64 ] || tar xf ${OCLCPUEXP_FN} -C oclcpuexp
[ -d fpgaemu/x64 ] || tar xf ${FPGAEMU_FN} -C fpgaemu
[ -d ${TBB_INSTALL_DIR}/lib ] || tar xf ${TBB_FN}
mkdir -p dpcpp_compiler/lib
mkdir -p dpcpp_compiler/lib/oclfpga
touch dpcpp_compiler/lib/oclfpga/fpgavars.sh
cp oclcpuexp/x64/libOpenCL.so* dpcpp_compiler/lib/
fi

- name: Install system components
Expand Down Expand Up @@ -121,7 +121,8 @@ jobs:
cat << 'EOF' > set_allvars.sh
#!/usr/bin/bash
export SYCL_BUNDLE_FOLDER=/home/runner/work/sycl_bundle
source ${SYCL_BUNDLE_FOLDER}/dpcpp_compiler/startup.sh
export PATH=${SYCL_BUNDLE_FOLDER}/dpcpp_compiler/bin:${PATH}
export LD_LIBRARY_PATH=${SYCL_BUNDLE_FOLDER}/dpcpp_compiler/lib:${LD_LIBRARY_PATH}
export LD_LIBRARY_PATH=${SYCL_BUNDLE_FOLDER}/oclcpuexp/x64:${LD_LIBRARY_PATH}
export LD_LIBRARY_PATH=${SYCL_BUNDLE_FOLDER}/fpgaemu/x64:${LD_LIBRARY_PATH}
export LD_LIBRARY_PATH=${SYCL_BUNDLE_FOLDER}/${TBB_INSTALL_DIR}/lib/intel64/gcc4.8:${LD_LIBRARY_PATH}
Expand Down
11 changes: 6 additions & 5 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -43,8 +43,8 @@ to program on XPUs.

# Installing

You can install the library with [conda](https://anaconda.org/intel/dpctl) and
[pip](https://pypi.org/project/dpctl/). It is also available in the [Intel(R)
You can install the library using [conda](https://anaconda.org/intel/dpctl) or
[pip](https://pypi.org/project/dpctl/) package managers. It is also available in the [Intel(R)
Distribution for
Python](https://www.intel.com/content/www/us/en/developer/tools/oneapi/distribution-for-python.html)
(IDP).
Expand All @@ -70,12 +70,13 @@ cloud, use the following command:
conda install dpctl -c intel
```

## PyPi
## Pip

To install `dpctl` from PyPi, run the following command:
The `dpctl` can be installed using `pip` obtaining wheel packages either from PyPi or from Intel(R) channel on Anaconda.
To install `dpctl` wheel package from Intel(R) channel on Anaconda, run the following command:

```bash
pip3 install dpctl
python -m pip install --index-url https://pypi.anaconda.org/intel/simple dpctl
```

Installing the bleeding edge
Expand Down
2 changes: 2 additions & 0 deletions conda-recipe/meta.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,8 @@ requirements:

test:
requires:
- {{ compiler('c') }}
- {{ compiler('cxx') }}
- cython
- setuptools
- pytest
Expand Down
7 changes: 6 additions & 1 deletion dpctl/tensor/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -58,10 +58,15 @@ set_source_files_properties(
${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/linear_sequences.cpp
${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/elementwise_functions.cpp
PROPERTIES COMPILE_OPTIONS "${_clang_prefix}-fno-fast-math")
if (UNIX)
set_source_files_properties(
${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/elementwise_functions.cpp
PROPERTIES COMPILE_DEFINITIONS "USE_STD_ABS_FOR_COMPLEX_TYPES;USE_STD_SQRT_FOR_COMPLEX_TYPES")
endif()
target_compile_options(${python_module_name} PRIVATE -fno-sycl-id-queries-fit-in-int)
target_link_options(${python_module_name} PRIVATE -fsycl-device-code-split=per_kernel)
if(UNIX)
# this option is support on Linux only
# this option is supported on Linux only
target_link_options(${python_module_name} PRIVATE -fsycl-link-huge-device-code)
endif()
target_include_directories(${python_module_name}
Expand Down
2 changes: 2 additions & 0 deletions dpctl/tensor/__init__.py
Original file line number Diff line number Diff line change
Expand Up @@ -158,6 +158,7 @@
trunc,
)
from ._reduction import sum
from ._testing import allclose

__all__ = [
"Device",
Expand Down Expand Up @@ -301,4 +302,5 @@
"tan",
"tanh",
"trunc",
"allclose",
]
85 changes: 80 additions & 5 deletions dpctl/tensor/_copy_utils.py
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
import dpctl.tensor as dpt
import dpctl.tensor._tensor_impl as ti
import dpctl.utils
from dpctl.tensor._ctors import _get_dtype
from dpctl.tensor._data_types import _get_dtype
from dpctl.tensor._device import normalize_queue_device

__doc__ = (
Expand Down Expand Up @@ -351,14 +351,16 @@ def _empty_like_orderK(X, dt, usm_type=None, dev=None):
)
st = list(X.strides)
perm = sorted(
range(X.ndim), key=lambda d: builtins.abs(st[d]), reverse=True
range(X.ndim),
key=lambda d: builtins.abs(st[d]) if X.shape[d] > 1 else 0,
reverse=True,
)
inv_perm = sorted(range(X.ndim), key=lambda i: perm[i])
st_sorted = [st[i] for i in perm]
sh = X.shape
sh_sorted = tuple(sh[i] for i in perm)
R = dpt.empty(sh_sorted, dtype=dt, usm_type=usm_type, device=dev, order="C")
if min(st_sorted) < 0:
if min(st) < 0:
st_sorted = [st[i] for i in perm]
sl = tuple(
slice(None, None, -1)
if st_sorted[i] < 0
Expand Down Expand Up @@ -395,9 +397,14 @@ def _empty_like_pair_orderK(X1, X2, dt, res_shape, usm_type, dev):
max_ndim = max(nd1, nd2)
st1 += [0] * (max_ndim - len(st1))
st2 += [0] * (max_ndim - len(st2))
sh1 = list(X1.shape) + [0] * (max_ndim - nd1)
sh2 = list(X2.shape) + [0] * (max_ndim - nd2)
perm = sorted(
range(max_ndim),
key=lambda d: (builtins.abs(st1[d]), builtins.abs(st2[d])),
key=lambda d: (
builtins.abs(st1[d]) if sh1[d] > 1 else 0,
builtins.abs(st2[d]) if sh2[d] > 1 else 0,
),
reverse=True,
)
inv_perm = sorted(range(max_ndim), key=lambda i: perm[i])
Expand All @@ -417,6 +424,74 @@ def _empty_like_pair_orderK(X1, X2, dt, res_shape, usm_type, dev):
return dpt.permute_dims(R, inv_perm)


def _empty_like_triple_orderK(X1, X2, X3, dt, res_shape, usm_type, dev):
if not isinstance(X1, dpt.usm_ndarray):
raise TypeError(f"Expected usm_ndarray, got {type(X1)}")
if not isinstance(X2, dpt.usm_ndarray):
raise TypeError(f"Expected usm_ndarray, got {type(X2)}")
if not isinstance(X3, dpt.usm_ndarray):
raise TypeError(f"Expected usm_ndarray, got {type(X3)}")
nd1 = X1.ndim
nd2 = X2.ndim
nd3 = X3.ndim
if X1.shape == res_shape and X2.shape == res_shape and len(res_shape) > nd3:
return _empty_like_pair_orderK(X1, X2, dt, res_shape, usm_type, dev)
elif (
X2.shape == res_shape and X3.shape == res_shape and len(res_shape) > nd1
):
return _empty_like_pair_orderK(X2, X3, dt, res_shape, usm_type, dev)
elif (
X1.shape == res_shape and X3.shape == res_shape and len(res_shape) > nd2
):
return _empty_like_pair_orderK(X1, X3, dt, res_shape, usm_type, dev)
fl1 = X1.flags
fl2 = X2.flags
fl3 = X3.flags
if fl1["C"] or fl2["C"] or fl3["C"]:
return dpt.empty(
res_shape, dtype=dt, usm_type=usm_type, device=dev, order="C"
)
if fl1["F"] and fl2["F"] and fl3["F"]:
return dpt.empty(
res_shape, dtype=dt, usm_type=usm_type, device=dev, order="F"
)
st1 = list(X1.strides)
st2 = list(X2.strides)
st3 = list(X3.strides)
max_ndim = max(nd1, nd2, nd3)
st1 += [0] * (max_ndim - len(st1))
st2 += [0] * (max_ndim - len(st2))
st3 += [0] * (max_ndim - len(st3))
sh1 = list(X1.shape) + [0] * (max_ndim - nd1)
sh2 = list(X2.shape) + [0] * (max_ndim - nd2)
sh3 = list(X3.shape) + [0] * (max_ndim - nd3)
perm = sorted(
range(max_ndim),
key=lambda d: (
builtins.abs(st1[d]) if sh1[d] > 1 else 0,
builtins.abs(st2[d]) if sh2[d] > 1 else 0,
builtins.abs(st3[d]) if sh3[d] > 1 else 0,
),
reverse=True,
)
inv_perm = sorted(range(max_ndim), key=lambda i: perm[i])
st1_sorted = [st1[i] for i in perm]
st2_sorted = [st2[i] for i in perm]
st3_sorted = [st3[i] for i in perm]
sh = res_shape
sh_sorted = tuple(sh[i] for i in perm)
R = dpt.empty(sh_sorted, dtype=dt, usm_type=usm_type, device=dev, order="C")
if max(min(st1_sorted), min(st2_sorted), min(st3_sorted)) < 0:
sl = tuple(
slice(None, None, -1)
if (st1_sorted[i] < 0 and st2_sorted[i] < 0 and st3_sorted[i] < 0)
else slice(None, None, None)
for i in range(nd1)
)
R = R[sl]
return dpt.permute_dims(R, inv_perm)


def copy(usm_ary, order="K"):
"""copy(ary, order="K")

Expand Down
43 changes: 3 additions & 40 deletions dpctl/tensor/_ctors.py
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,8 @@
import dpctl.tensor as dpt
import dpctl.tensor._tensor_impl as ti
import dpctl.utils
from dpctl.tensor._copy_utils import _empty_like_orderK
from dpctl.tensor._data_types import _get_dtype
from dpctl.tensor._device import normalize_queue_device
from dpctl.tensor._usmarray import _is_object_with_buffer_protocol

Expand All @@ -32,24 +34,6 @@
_host_set = frozenset([None])


def _get_dtype(dtype, sycl_obj, ref_type=None):
if dtype is None:
if ref_type in [None, float] or np.issubdtype(ref_type, np.floating):
dtype = ti.default_device_fp_type(sycl_obj)
return dpt.dtype(dtype)
if ref_type in [bool, np.bool_]:
dtype = ti.default_device_bool_type(sycl_obj)
return dpt.dtype(dtype)
if ref_type is int or np.issubdtype(ref_type, np.integer):
dtype = ti.default_device_int_type(sycl_obj)
return dpt.dtype(dtype)
if ref_type is complex or np.issubdtype(ref_type, np.complexfloating):
dtype = ti.default_device_complex_type(sycl_obj)
return dpt.dtype(dtype)
raise TypeError(f"Reference type {ref_type} not recognized.")
return dpt.dtype(dtype)


def _array_info_dispatch(obj):
if isinstance(obj, dpt.usm_ndarray):
return obj.shape, obj.dtype, frozenset([obj.sycl_queue])
Expand Down Expand Up @@ -162,28 +146,7 @@ def _asarray_from_usm_ndarray(
order = "C" if c_contig else "F"
if order == "K":
_ensure_native_dtype_device_support(dtype, copy_q.sycl_device)
# new USM allocation
res = dpt.usm_ndarray(
usm_ndary.shape,
dtype=dtype,
buffer=usm_type,
order="C",
buffer_ctor_kwargs={"queue": copy_q},
)
original_strides = usm_ndary.strides
ind = sorted(
range(usm_ndary.ndim),
key=lambda i: abs(original_strides[i]),
reverse=True,
)
new_strides = tuple(res.strides[ind[i]] for i in ind)
# reuse previously made USM allocation
res = dpt.usm_ndarray(
usm_ndary.shape,
dtype=res.dtype,
buffer=res.usm_data,
strides=new_strides,
)
res = _empty_like_orderK(usm_ndary, dtype, usm_type, copy_q)
else:
_ensure_native_dtype_device_support(dtype, copy_q.sycl_device)
res = dpt.usm_ndarray(
Expand Down
44 changes: 44 additions & 0 deletions dpctl/tensor/_data_types.py
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,25 @@
# See the License for the specific language governing permissions and
# limitations under the License.

from numpy import bool_ as np_bool_
from numpy import complexfloating as np_complexfloating
from numpy import dtype
from numpy import floating as np_floating
from numpy import integer as np_integer
from numpy import issubdtype as np_issubdtype

from dpctl.tensor._tensor_impl import (
default_device_bool_type as ti_default_device_bool_type,
)
from dpctl.tensor._tensor_impl import (
default_device_complex_type as ti_default_device_complex_type,
)
from dpctl.tensor._tensor_impl import (
default_device_fp_type as ti_default_device_fp_type,
)
from dpctl.tensor._tensor_impl import (
default_device_int_type as ti_default_device_int_type,
)

bool = dtype("bool")
int8 = dtype("int8")
Expand Down Expand Up @@ -74,6 +92,32 @@ def isdtype(dtype_, kind):
raise TypeError(f"Unsupported data type kind: {kind}")


def _get_dtype(inp_dt, sycl_obj, ref_type=None):
"""
Type inference utility to construct data type
object with defaults based on reference type.

_get_dtype is used by dpctl.tensor.asarray
to infer data type of the output array from the
input sequence.
"""
if inp_dt is None:
if ref_type in [None, float] or np_issubdtype(ref_type, np_floating):
fp_dt = ti_default_device_fp_type(sycl_obj)
return dtype(fp_dt)
if ref_type in [bool, np_bool_]:
bool_dt = ti_default_device_bool_type(sycl_obj)
return dtype(bool_dt)
if ref_type is int or np_issubdtype(ref_type, np_integer):
int_dt = ti_default_device_int_type(sycl_obj)
return dtype(int_dt)
if ref_type is complex or np_issubdtype(ref_type, np_complexfloating):
cfp_dt = ti_default_device_complex_type(sycl_obj)
return dtype(cfp_dt)
raise TypeError(f"Reference type {ref_type} not recognized.")
return dtype(inp_dt)


__all__ = [
"dtype",
"isdtype",
Expand Down
Loading