diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml index f6dc0dcc72..f5496bae39 100644 --- a/.github/workflows/conda-package.yml +++ b/.github/workflows/conda-package.yml @@ -30,7 +30,7 @@ jobs: fail-fast: false matrix: python: ['3.9', '3.10', '3.11'] - os: [ubuntu-latest, windows-latest] + os: [ubuntu-latest, windows-2019] runs-on: ${{ matrix.os }} diff --git a/.github/workflows/coverage.yml b/.github/workflows/coverage.yml index c949f6dfc6..c9b2bc9154 100644 --- a/.github/workflows/coverage.yml +++ b/.github/workflows/coverage.yml @@ -54,7 +54,7 @@ jobs: # Ignoring test due to opencl driver optimization bug - name: Run tests with coverage run: | - pytest -q --cov --cov-report xml --pyargs numba_dpex \ + pytest -q --cov=./ --cov-report xml --pyargs numba_dpex \ -k 'not test_1d_strided_dpnp_array_in_kernel[2]' - name: Coveralls diff --git a/.github/workflows/cpp_style_checks.yml b/.github/workflows/cpp_style_checks.yml index 71c431b760..92ead9fb16 100644 --- a/.github/workflows/cpp_style_checks.yml +++ b/.github/workflows/cpp_style_checks.yml @@ -18,7 +18,7 @@ jobs: steps: - uses: actions/checkout@v4 - name: Run clang-format style check for C/C++ programs. - uses: jidicula/clang-format-action@v4.11.0 + uses: jidicula/clang-format-action@v4.13.0 with: clang-format-version: '14' check-path: 'numba_dpex/dpctl_iface' diff --git a/.github/workflows/openssf-scorecard.yml b/.github/workflows/openssf-scorecard.yml index 04a32f7f54..6f953e3b8d 100644 --- a/.github/workflows/openssf-scorecard.yml +++ b/.github/workflows/openssf-scorecard.yml @@ -40,7 +40,7 @@ jobs: persist-credentials: false - name: "Run analysis" - uses: ossf/scorecard-action@0864cf19026789058feabb7e87baa5f140aac736 # v2.3.1 + uses: ossf/scorecard-action@dc50aa9510b46c811795eb24b2f1ba02a914e534 # v2.3.3 with: results_file: results.sarif results_format: sarif @@ -71,6 +71,6 @@ jobs: # Upload the results to GitHub's code scanning dashboard. - name: "Upload to code-scanning" if: ${{ github.event_name == 'push' }} - uses: github/codeql-action/upload-sarif@d39d31e687223d841ef683f52467bd88e9b21c14 # v3.25.3 + uses: github/codeql-action/upload-sarif@9fdb3e49720b44c48891d036bb502feb25684276 # v3.25.6 with: sarif_file: results.sarif diff --git a/.gitignore b/.gitignore index bd383c0a5b..4735becd1e 100644 --- a/.gitignore +++ b/.gitignore @@ -26,6 +26,9 @@ _skbuild docs/source/developer/autogen* +# Ignore versioneer generated files +numba_dpex/_version.py + # Ignore generated cpp files numba_dpex/dpnp_iface/*.cpp numba_dpex/dpnp_iface/*.h diff --git a/CHANGELOG.md b/CHANGELOG.md index ebab23aa0c..56f5f6bee8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -4,13 +4,21 @@ All notable changes to this project will be documented in this file. The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.0.0/), and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html). -## [0.23.0] - 2024-04-XX +## [0.23.0] - 2024-05-28 ### Fixed * Array alignment problem for stack arrays allocated for kernel arguments. (#1357) * Issue #892, #906 caused by incorrect code generation for indexing (#1377) -* Fix `KernelHasReturnValueError` inside `KernelDispatcher`. (#1394) +* Generation of `KernelHasReturnValueError` error inside `KernelDispatcher`. (#1394) * Issue #1390: broken support for slicing into `dpctl.tensor.usm_ndarray` in kernels (#1425) +* Support for Wheels package on Windows (#1430) +* Incorrect mangled name for kernel function arguments (#1443) +* Remove artifacts from conda/wheel packages residing in root level (#1450) +* GDB tests to work properly on Intel Max GPU (#1451) +* Improper wheels installation on unsupported platforms (#1452) +* Ref-counting of Python object temporaries in unboxing code (#1454) +* Segfault caused by using `malloc` to allocate `NRT_MemInfo`. Replaced with Numba's NRT `alloc` (#1458) +* Incorrect package name in README.md (#1463) ### Added * A new overloaded `dimensions` attribute for all index-space id classes (#1359) @@ -24,18 +32,24 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * A `sycl::local_accessor`-like API (`kernel_api.LocalAccessor`) for numba-dpex kernel (#1331) * Specialization support for `device_func` decorator (#1398) * Support for all `kernel_api` functions inside the `numba_dpex.kernel` decorator. (#1400) +* Support for dpnp 0.15 (#1434, #1464) +* Improvements to pyproject.toml configs to build numba-dpex from source. (#1449) +* Load the `SPV_INTEL_variable_length_array` SPIR-V extension to supporting arrays in private address-space on Intel Max GPU. (#1451) ### Changed * Default inline threshold value set to `2` from `None`. (#1385) * Port parfor kernel templates to `kernel_api` (#1416), (#1424) +* Use `SPIRVKernelDispatcher` for parfor kernel dispatch (#1435, #1448) +* All examples use the latest dpctl API (#1431) * Minimum required dpctl version is now 0.16.1 -* Minimum required numba version is now 0.59.0 +* Minimum required numba version is now 0.59.0 (#1462) ### Removed * OpenCL-like kernel API functions (#1420) * `func` decorator (replaced by `device_func`) (#1400) * `numba_dpex.experimental.kernel` and `numba_dpex.experimental.device_func` (#1400) + ## [0.22.0] - 2024-02-19 ### Fixed diff --git a/MANIFEST.in b/MANIFEST.in deleted file mode 100644 index 8c457d17ad..0000000000 --- a/MANIFEST.in +++ /dev/null @@ -1,6 +0,0 @@ -include MANIFEST.in -include README.md setup.py LICENSE - -include numba_dpex/_version.py - -recursive-include numba_dpex/examples * diff --git a/README.md b/README.md index dcd5f51638..84a1afde6d 100644 --- a/README.md +++ b/README.md @@ -12,7 +12,7 @@ -Data-parallel Extension for Numba* (numba-dpex) is an open-source standalone +Data Parallel Extension for Numba* (numba-dpex) is an open-source standalone extension for the [Numba](http://numba.pydata.org) Python JIT compiler. Numba-dpex provides a [SYCL*](https://sycl.tech/)-like API for kernel programming Python. SYCL* is an open standard developed by the [Unified diff --git a/conda-recipe/bld.bat b/conda-recipe/bld.bat index 1a74face5c..277662827a 100644 --- a/conda-recipe/bld.bat +++ b/conda-recipe/bld.bat @@ -35,7 +35,13 @@ for /f %%f in ('dir /b /S .\dist') do ( :: wheel file was renamed for /f %%f in ('dir /b /S .\dist') do ( - %PYTHON% -m pip install %%f + %PYTHON% -m pip install %%f ^ + --no-build-isolation ^ + --no-deps ^ + --only-binary :all: ^ + --no-index ^ + --prefix %PREFIX% ^ + -vv if %ERRORLEVEL% neq 0 exit 1 ) diff --git a/conda-recipe/build.sh b/conda-recipe/build.sh index 68f12528f9..a6fb6607f7 100644 --- a/conda-recipe/build.sh +++ b/conda-recipe/build.sh @@ -9,6 +9,9 @@ echo "--gcc-toolchain=${BUILD_PREFIX} --sysroot=${BUILD_PREFIX}/${HOST}/sysroot ICPXCFG="$(pwd)/icpx_for_conda.cfg" ICXCFG="$(pwd)/icpx_for_conda.cfg" +read -r GLIBC_MAJOR GLIBC_MINOR <<<"$(conda list '^sysroot_linux-64$' \ + | tail -n 1 | awk '{print $2}' | grep -oP '\d+' | head -n 2 | tr '\n' ' ')" + export ICXCFG export ICPXCFG @@ -26,8 +29,15 @@ export PATH=$CONDA_PREFIX/bin-llvm:$PATH # -wnx flags mean: --wheel --no-isolation --skip-dependency-check ${PYTHON} -m build -w -n -x ${PYTHON} -m wheel tags --remove --build "$GIT_DESCRIBE_NUMBER" \ - --platform-tag manylinux2014_x86_64 dist/numba_dpex*.whl -${PYTHON} -m pip install dist/numba_dpex*.whl + --platform-tag "manylinux_${GLIBC_MAJOR}_${GLIBC_MINOR}_x86_64" \ + dist/numba_dpex*.whl +${PYTHON} -m pip install dist/numba_dpex*.whl \ + --no-build-isolation \ + --no-deps \ + --only-binary :all: \ + --no-index \ + --prefix "${PREFIX}" \ + -vv # Copy wheel package if [[ -v WHEELS_OUTPUT_FOLDER ]]; then diff --git a/conda-recipe/meta.yaml b/conda-recipe/meta.yaml index 8db61810fc..6e27f82491 100644 --- a/conda-recipe/meta.yaml +++ b/conda-recipe/meta.yaml @@ -3,6 +3,9 @@ {% set excluded_compiler_version2 = "2024.0.2" %} {% set excluded_compiler_version3 = "2024.0.3" %} +{% set pyproject = load_file_data('pyproject.toml') %} +{% set py_build_deps = pyproject.get('build-system', {}).get('requires', []) %} + package: name: numba-dpex version: {{ GIT_DESCRIBE_TAG }} @@ -21,38 +24,39 @@ requirements: - {{ compiler('cxx') }} - {{ compiler('dpcpp') }} >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }},!={{ excluded_compiler_version3 }} # [win] - {{ compiler('dpcpp') }} >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }} # [linux] - # specific version of sysroot required by dpcpp, but 2024.0.0 package - # does not have it in meta data - - sysroot_linux-64 >=2.28 # [linux] + # Minimal supported version of sysroot (which is version of glibc) to + # have compatibility with wider range of linux distributions. + # 2.28 is the minimal supported version by dpcpp + - sysroot_linux-64 =2.28 # [linux] host: + - python + - pip >=24.0 - dpcpp-cpp-rt >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }},!={{ excluded_compiler_version3 }} # [win] - dpcpp-cpp-rt >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }} # [linux] - - python + # ensure we are using latest version of setuptools, since we don't need + # editable environments for release. - setuptools >=69 - - scikit-build >=0.17 - - ninja >=1.11.1 # [not win] - - cmake >=3.29 - - numba >=0.59 - - llvmlite >=0.42.0 - - dpctl >=0.16.1 - - dpnp >=0.14 - - numpy >=1.24 - # TODO: temporary fix, because IGC is broken for output produced by llvm-spirv 2024.1 on windows - - dpcpp-llvm-spirv >={{ required_compiler_version }} # [not win] - - dpcpp-llvm-spirv >={{ required_compiler_version }},<2024.1 # [win] - - wheel >=0.43 - - pip >=24.0 - - python-build >=1.1 - - versioneer==0.29 + {% for dep in py_build_deps %} + {% if dep.startswith('ninja') %} + - {{ dep.split(';')[0] }} # [not win] + {% elif dep.startswith('cmake') %} + - {{ dep }} + {% elif dep.startswith('build>=') %} + - {{ 'python-' ~ dep }} + {% else %} + - {{ dep|replace('_','-') }} + {% endif %} + {% endfor %} # versioneer dependency - tomli # [py<311] + # While we don't need it for build, but install it here, so we can + # pin_compatible at run section. + - dpcpp-llvm-spirv >={{ required_compiler_version }} run: - {{ pin_compatible('dpcpp-cpp-rt', min_pin='x.x', max_pin='x') }} - {{ pin_compatible('intel-cmplr-lib-rt', min_pin='x.x', max_pin='x') }} + # TODO: pick up min version from dep - {{ pin_compatible('dpcpp-llvm-spirv', min_pin='x.x', max_pin='x') }} - # TODO: temporary fix, because IGC is broken for output produced by llvm-spirv 2024.1 on windows - - {{ pin_compatible('dpcpp-llvm-spirv', min_pin='x.x', max_pin='x') }} # [not win] - - {{ pin_compatible('dpcpp-llvm-spirv', min_pin='x.x', max_pin='x', upper_bound='2024.1') }} # [win] - {{ pin_compatible('dpnp', min_pin='x.x.x', max_pin='x.x') }} - {{ pin_compatible('dpctl', min_pin='x.x.x', max_pin='x.x') }} - {{ pin_compatible('numba', min_pin='x.x.x', max_pin='x.x') }} diff --git a/numba_dpex/__init__.py b/numba_dpex/__init__.py index 4a24e28502..52228ac303 100644 --- a/numba_dpex/__init__.py +++ b/numba_dpex/__init__.py @@ -17,7 +17,7 @@ from numba import __version__ as numba_version from .kernel_api_impl.spirv import target as spirv_kernel_target -from .numba_patches import patch_is_ufunc +from .numba_patches import patch_ufuncs from .register_kernel_api_overloads import init_kernel_api_spirv_overloads @@ -70,7 +70,7 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]: dpctl_sem_version = parse_sem_version(dpctl.__version__) # Monkey patches -patch_is_ufunc.patch() +patch_ufuncs.patch() from numba import prange # noqa E402 diff --git a/numba_dpex/_version.py b/numba_dpex/_version.py deleted file mode 100644 index f076bfbd66..0000000000 --- a/numba_dpex/_version.py +++ /dev/null @@ -1,683 +0,0 @@ - -# This file helps to compute a version number in source trees obtained from -# git-archive tarball (such as those provided by githubs download-from-tag -# feature). Distribution tarballs (built by setup.py sdist) and build -# directories (produced by setup.py build) will contain a much shorter file -# that just contains the computed version number. - -# This file is released into the public domain. -# Generated by versioneer-0.29 -# https://github.com/python-versioneer/python-versioneer - -"""Git implementation of _version.py.""" - -import errno -import os -import re -import subprocess -import sys -from typing import Any, Callable, Dict, List, Optional, Tuple -import functools - - -def get_keywords() -> Dict[str, str]: - """Get the keywords needed to look up the version information.""" - # these strings will be replaced by git during git-archive. - # setup.py/versioneer.py will grep for the variable names, so they must - # each be defined on a line of their own. _version.py will just call - # get_keywords(). - git_refnames = "$Format:%d$" - git_full = "$Format:%H$" - git_date = "$Format:%ci$" - keywords = {"refnames": git_refnames, "full": git_full, "date": git_date} - return keywords - - -class VersioneerConfig: - """Container for Versioneer configuration parameters.""" - - VCS: str - style: str - tag_prefix: str - parentdir_prefix: str - versionfile_source: str - verbose: bool - - -def get_config() -> VersioneerConfig: - """Create, populate and return the VersioneerConfig() object.""" - # these strings are filled in when 'setup.py versioneer' creates - # _version.py - cfg = VersioneerConfig() - cfg.VCS = "git" - cfg.style = "pep440" - cfg.tag_prefix = "" - cfg.parentdir_prefix = "" - cfg.versionfile_source = "numba_dpex/_version.py" - cfg.verbose = False - return cfg - - -class NotThisMethod(Exception): - """Exception raised if a method is not valid for the current scenario.""" - - -LONG_VERSION_PY: Dict[str, str] = {} -HANDLERS: Dict[str, Dict[str, Callable]] = {} - - -def register_vcs_handler(vcs: str, method: str) -> Callable: # decorator - """Create decorator to mark a method as the handler of a VCS.""" - def decorate(f: Callable) -> Callable: - """Store f in HANDLERS[vcs][method].""" - if vcs not in HANDLERS: - HANDLERS[vcs] = {} - HANDLERS[vcs][method] = f - return f - return decorate - - -def run_command( - commands: List[str], - args: List[str], - cwd: Optional[str] = None, - verbose: bool = False, - hide_stderr: bool = False, - env: Optional[Dict[str, str]] = None, -) -> Tuple[Optional[str], Optional[int]]: - """Call the given command(s).""" - assert isinstance(commands, list) - process = None - - popen_kwargs: Dict[str, Any] = {} - if sys.platform == "win32": - # This hides the console window if pythonw.exe is used - startupinfo = subprocess.STARTUPINFO() - startupinfo.dwFlags |= subprocess.STARTF_USESHOWWINDOW - popen_kwargs["startupinfo"] = startupinfo - - for command in commands: - try: - dispcmd = str([command] + args) - # remember shell=False, so use git.cmd on windows, not just git - process = subprocess.Popen([command] + args, cwd=cwd, env=env, - stdout=subprocess.PIPE, - stderr=(subprocess.PIPE if hide_stderr - else None), **popen_kwargs) - break - except OSError as e: - if e.errno == errno.ENOENT: - continue - if verbose: - print("unable to run %s" % dispcmd) - print(e) - return None, None - else: - if verbose: - print("unable to find command, tried %s" % (commands,)) - return None, None - stdout = process.communicate()[0].strip().decode() - if process.returncode != 0: - if verbose: - print("unable to run %s (error)" % dispcmd) - print("stdout was %s" % stdout) - return None, process.returncode - return stdout, process.returncode - - -def versions_from_parentdir( - parentdir_prefix: str, - root: str, - verbose: bool, -) -> Dict[str, Any]: - """Try to determine the version from the parent directory name. - - Source tarballs conventionally unpack into a directory that includes both - the project name and a version string. We will also support searching up - two directory levels for an appropriately named parent directory - """ - rootdirs = [] - - for _ in range(3): - dirname = os.path.basename(root) - if dirname.startswith(parentdir_prefix): - return {"version": dirname[len(parentdir_prefix):], - "full-revisionid": None, - "dirty": False, "error": None, "date": None} - rootdirs.append(root) - root = os.path.dirname(root) # up a level - - if verbose: - print("Tried directories %s but none started with prefix %s" % - (str(rootdirs), parentdir_prefix)) - raise NotThisMethod("rootdir doesn't start with parentdir_prefix") - - -@register_vcs_handler("git", "get_keywords") -def git_get_keywords(versionfile_abs: str) -> Dict[str, str]: - """Extract version information from the given file.""" - # the code embedded in _version.py can just fetch the value of these - # keywords. When used from setup.py, we don't want to import _version.py, - # so we do it with a regexp instead. This function is not used from - # _version.py. - keywords: Dict[str, str] = {} - try: - with open(versionfile_abs, "r") as fobj: - for line in fobj: - if line.strip().startswith("git_refnames ="): - mo = re.search(r'=\s*"(.*)"', line) - if mo: - keywords["refnames"] = mo.group(1) - if line.strip().startswith("git_full ="): - mo = re.search(r'=\s*"(.*)"', line) - if mo: - keywords["full"] = mo.group(1) - if line.strip().startswith("git_date ="): - mo = re.search(r'=\s*"(.*)"', line) - if mo: - keywords["date"] = mo.group(1) - except OSError: - pass - return keywords - - -@register_vcs_handler("git", "keywords") -def git_versions_from_keywords( - keywords: Dict[str, str], - tag_prefix: str, - verbose: bool, -) -> Dict[str, Any]: - """Get version information from git keywords.""" - if "refnames" not in keywords: - raise NotThisMethod("Short version file found") - date = keywords.get("date") - if date is not None: - # Use only the last line. Previous lines may contain GPG signature - # information. - date = date.splitlines()[-1] - - # git-2.2.0 added "%cI", which expands to an ISO-8601 -compliant - # datestamp. However we prefer "%ci" (which expands to an "ISO-8601 - # -like" string, which we must then edit to make compliant), because - # it's been around since git-1.5.3, and it's too difficult to - # discover which version we're using, or to work around using an - # older one. - date = date.strip().replace(" ", "T", 1).replace(" ", "", 1) - refnames = keywords["refnames"].strip() - if refnames.startswith("$Format"): - if verbose: - print("keywords are unexpanded, not using") - raise NotThisMethod("unexpanded keywords, not a git-archive tarball") - refs = {r.strip() for r in refnames.strip("()").split(",")} - # starting in git-1.8.3, tags are listed as "tag: foo-1.0" instead of - # just "foo-1.0". If we see a "tag: " prefix, prefer those. - TAG = "tag: " - tags = {r[len(TAG):] for r in refs if r.startswith(TAG)} - if not tags: - # Either we're using git < 1.8.3, or there really are no tags. We use - # a heuristic: assume all version tags have a digit. The old git %d - # expansion behaves like git log --decorate=short and strips out the - # refs/heads/ and refs/tags/ prefixes that would let us distinguish - # between branches and tags. By ignoring refnames without digits, we - # filter out many common branch names like "release" and - # "stabilization", as well as "HEAD" and "master". - tags = {r for r in refs if re.search(r'\d', r)} - if verbose: - print("discarding '%s', no digits" % ",".join(refs - tags)) - if verbose: - print("likely tags: %s" % ",".join(sorted(tags))) - for ref in sorted(tags): - # sorting will prefer e.g. "2.0" over "2.0rc1" - if ref.startswith(tag_prefix): - r = ref[len(tag_prefix):] - # Filter out refs that exactly match prefix or that don't start - # with a number once the prefix is stripped (mostly a concern - # when prefix is '') - if not re.match(r'\d', r): - continue - if verbose: - print("picking %s" % r) - return {"version": r, - "full-revisionid": keywords["full"].strip(), - "dirty": False, "error": None, - "date": date} - # no suitable tags, so version is "0+unknown", but full hex is still there - if verbose: - print("no suitable tags, using unknown + full revision id") - return {"version": "0+unknown", - "full-revisionid": keywords["full"].strip(), - "dirty": False, "error": "no suitable tags", "date": None} - - -@register_vcs_handler("git", "pieces_from_vcs") -def git_pieces_from_vcs( - tag_prefix: str, - root: str, - verbose: bool, - runner: Callable = run_command -) -> Dict[str, Any]: - """Get version from 'git describe' in the root of the source tree. - - This only gets called if the git-archive 'subst' keywords were *not* - expanded, and _version.py hasn't already been rewritten with a short - version string, meaning we're inside a checked out source tree. - """ - GITS = ["git"] - if sys.platform == "win32": - GITS = ["git.cmd", "git.exe"] - - # GIT_DIR can interfere with correct operation of Versioneer. - # It may be intended to be passed to the Versioneer-versioned project, - # but that should not change where we get our version from. - env = os.environ.copy() - env.pop("GIT_DIR", None) - runner = functools.partial(runner, env=env) - - _, rc = runner(GITS, ["rev-parse", "--git-dir"], cwd=root, - hide_stderr=not verbose) - if rc != 0: - if verbose: - print("Directory %s not under git control" % root) - raise NotThisMethod("'git rev-parse --git-dir' returned error") - - # if there is a tag matching tag_prefix, this yields TAG-NUM-gHEX[-dirty] - # if there isn't one, this yields HEX[-dirty] (no NUM) - describe_out, rc = runner(GITS, [ - "describe", "--tags", "--dirty", "--always", "--long", - "--match", f"{tag_prefix}[[:digit:]]*" - ], cwd=root) - # --long was added in git-1.5.5 - if describe_out is None: - raise NotThisMethod("'git describe' failed") - describe_out = describe_out.strip() - full_out, rc = runner(GITS, ["rev-parse", "HEAD"], cwd=root) - if full_out is None: - raise NotThisMethod("'git rev-parse' failed") - full_out = full_out.strip() - - pieces: Dict[str, Any] = {} - pieces["long"] = full_out - pieces["short"] = full_out[:7] # maybe improved later - pieces["error"] = None - - branch_name, rc = runner(GITS, ["rev-parse", "--abbrev-ref", "HEAD"], - cwd=root) - # --abbrev-ref was added in git-1.6.3 - if rc != 0 or branch_name is None: - raise NotThisMethod("'git rev-parse --abbrev-ref' returned error") - branch_name = branch_name.strip() - - if branch_name == "HEAD": - # If we aren't exactly on a branch, pick a branch which represents - # the current commit. If all else fails, we are on a branchless - # commit. - branches, rc = runner(GITS, ["branch", "--contains"], cwd=root) - # --contains was added in git-1.5.4 - if rc != 0 or branches is None: - raise NotThisMethod("'git branch --contains' returned error") - branches = branches.split("\n") - - # Remove the first line if we're running detached - if "(" in branches[0]: - branches.pop(0) - - # Strip off the leading "* " from the list of branches. - branches = [branch[2:] for branch in branches] - if "master" in branches: - branch_name = "master" - elif not branches: - branch_name = None - else: - # Pick the first branch that is returned. Good or bad. - branch_name = branches[0] - - pieces["branch"] = branch_name - - # parse describe_out. It will be like TAG-NUM-gHEX[-dirty] or HEX[-dirty] - # TAG might have hyphens. - git_describe = describe_out - - # look for -dirty suffix - dirty = git_describe.endswith("-dirty") - pieces["dirty"] = dirty - if dirty: - git_describe = git_describe[:git_describe.rindex("-dirty")] - - # now we have TAG-NUM-gHEX or HEX - - if "-" in git_describe: - # TAG-NUM-gHEX - mo = re.search(r'^(.+)-(\d+)-g([0-9a-f]+)$', git_describe) - if not mo: - # unparsable. Maybe git-describe is misbehaving? - pieces["error"] = ("unable to parse git-describe output: '%s'" - % describe_out) - return pieces - - # tag - full_tag = mo.group(1) - if not full_tag.startswith(tag_prefix): - if verbose: - fmt = "tag '%s' doesn't start with prefix '%s'" - print(fmt % (full_tag, tag_prefix)) - pieces["error"] = ("tag '%s' doesn't start with prefix '%s'" - % (full_tag, tag_prefix)) - return pieces - pieces["closest-tag"] = full_tag[len(tag_prefix):] - - # distance: number of commits since tag - pieces["distance"] = int(mo.group(2)) - - # commit: short hex revision ID - pieces["short"] = mo.group(3) - - else: - # HEX: no tags - pieces["closest-tag"] = None - out, rc = runner(GITS, ["rev-list", "HEAD", "--left-right"], cwd=root) - pieces["distance"] = len(out.split()) # total number of commits - - # commit date: see ISO-8601 comment in git_versions_from_keywords() - date = runner(GITS, ["show", "-s", "--format=%ci", "HEAD"], cwd=root)[0].strip() - # Use only the last line. Previous lines may contain GPG signature - # information. - date = date.splitlines()[-1] - pieces["date"] = date.strip().replace(" ", "T", 1).replace(" ", "", 1) - - return pieces - - -def plus_or_dot(pieces: Dict[str, Any]) -> str: - """Return a + if we don't already have one, else return a .""" - if "+" in pieces.get("closest-tag", ""): - return "." - return "+" - - -def render_pep440(pieces: Dict[str, Any]) -> str: - """Build up version string, with post-release "local version identifier". - - Our goal: TAG[+DISTANCE.gHEX[.dirty]] . Note that if you - get a tagged build and then dirty it, you'll get TAG+0.gHEX.dirty - - Exceptions: - 1: no tags. git_describe was just HEX. 0+untagged.DISTANCE.gHEX[.dirty] - """ - if pieces["closest-tag"]: - rendered = pieces["closest-tag"] - if pieces["distance"] or pieces["dirty"]: - rendered += plus_or_dot(pieces) - rendered += "%d.g%s" % (pieces["distance"], pieces["short"]) - if pieces["dirty"]: - rendered += ".dirty" - else: - # exception #1 - rendered = "0+untagged.%d.g%s" % (pieces["distance"], - pieces["short"]) - if pieces["dirty"]: - rendered += ".dirty" - return rendered - - -def render_pep440_branch(pieces: Dict[str, Any]) -> str: - """TAG[[.dev0]+DISTANCE.gHEX[.dirty]] . - - The ".dev0" means not master branch. Note that .dev0 sorts backwards - (a feature branch will appear "older" than the master branch). - - Exceptions: - 1: no tags. 0[.dev0]+untagged.DISTANCE.gHEX[.dirty] - """ - if pieces["closest-tag"]: - rendered = pieces["closest-tag"] - if pieces["distance"] or pieces["dirty"]: - if pieces["branch"] != "master": - rendered += ".dev0" - rendered += plus_or_dot(pieces) - rendered += "%d.g%s" % (pieces["distance"], pieces["short"]) - if pieces["dirty"]: - rendered += ".dirty" - else: - # exception #1 - rendered = "0" - if pieces["branch"] != "master": - rendered += ".dev0" - rendered += "+untagged.%d.g%s" % (pieces["distance"], - pieces["short"]) - if pieces["dirty"]: - rendered += ".dirty" - return rendered - - -def pep440_split_post(ver: str) -> Tuple[str, Optional[int]]: - """Split pep440 version string at the post-release segment. - - Returns the release segments before the post-release and the - post-release version number (or -1 if no post-release segment is present). - """ - vc = str.split(ver, ".post") - return vc[0], int(vc[1] or 0) if len(vc) == 2 else None - - -def render_pep440_pre(pieces: Dict[str, Any]) -> str: - """TAG[.postN.devDISTANCE] -- No -dirty. - - Exceptions: - 1: no tags. 0.post0.devDISTANCE - """ - if pieces["closest-tag"]: - if pieces["distance"]: - # update the post release segment - tag_version, post_version = pep440_split_post(pieces["closest-tag"]) - rendered = tag_version - if post_version is not None: - rendered += ".post%d.dev%d" % (post_version + 1, pieces["distance"]) - else: - rendered += ".post0.dev%d" % (pieces["distance"]) - else: - # no commits, use the tag as the version - rendered = pieces["closest-tag"] - else: - # exception #1 - rendered = "0.post0.dev%d" % pieces["distance"] - return rendered - - -def render_pep440_post(pieces: Dict[str, Any]) -> str: - """TAG[.postDISTANCE[.dev0]+gHEX] . - - The ".dev0" means dirty. Note that .dev0 sorts backwards - (a dirty tree will appear "older" than the corresponding clean one), - but you shouldn't be releasing software with -dirty anyways. - - Exceptions: - 1: no tags. 0.postDISTANCE[.dev0] - """ - if pieces["closest-tag"]: - rendered = pieces["closest-tag"] - if pieces["distance"] or pieces["dirty"]: - rendered += ".post%d" % pieces["distance"] - if pieces["dirty"]: - rendered += ".dev0" - rendered += plus_or_dot(pieces) - rendered += "g%s" % pieces["short"] - else: - # exception #1 - rendered = "0.post%d" % pieces["distance"] - if pieces["dirty"]: - rendered += ".dev0" - rendered += "+g%s" % pieces["short"] - return rendered - - -def render_pep440_post_branch(pieces: Dict[str, Any]) -> str: - """TAG[.postDISTANCE[.dev0]+gHEX[.dirty]] . - - The ".dev0" means not master branch. - - Exceptions: - 1: no tags. 0.postDISTANCE[.dev0]+gHEX[.dirty] - """ - if pieces["closest-tag"]: - rendered = pieces["closest-tag"] - if pieces["distance"] or pieces["dirty"]: - rendered += ".post%d" % pieces["distance"] - if pieces["branch"] != "master": - rendered += ".dev0" - rendered += plus_or_dot(pieces) - rendered += "g%s" % pieces["short"] - if pieces["dirty"]: - rendered += ".dirty" - else: - # exception #1 - rendered = "0.post%d" % pieces["distance"] - if pieces["branch"] != "master": - rendered += ".dev0" - rendered += "+g%s" % pieces["short"] - if pieces["dirty"]: - rendered += ".dirty" - return rendered - - -def render_pep440_old(pieces: Dict[str, Any]) -> str: - """TAG[.postDISTANCE[.dev0]] . - - The ".dev0" means dirty. - - Exceptions: - 1: no tags. 0.postDISTANCE[.dev0] - """ - if pieces["closest-tag"]: - rendered = pieces["closest-tag"] - if pieces["distance"] or pieces["dirty"]: - rendered += ".post%d" % pieces["distance"] - if pieces["dirty"]: - rendered += ".dev0" - else: - # exception #1 - rendered = "0.post%d" % pieces["distance"] - if pieces["dirty"]: - rendered += ".dev0" - return rendered - - -def render_git_describe(pieces: Dict[str, Any]) -> str: - """TAG[-DISTANCE-gHEX][-dirty]. - - Like 'git describe --tags --dirty --always'. - - Exceptions: - 1: no tags. HEX[-dirty] (note: no 'g' prefix) - """ - if pieces["closest-tag"]: - rendered = pieces["closest-tag"] - if pieces["distance"]: - rendered += "-%d-g%s" % (pieces["distance"], pieces["short"]) - else: - # exception #1 - rendered = pieces["short"] - if pieces["dirty"]: - rendered += "-dirty" - return rendered - - -def render_git_describe_long(pieces: Dict[str, Any]) -> str: - """TAG-DISTANCE-gHEX[-dirty]. - - Like 'git describe --tags --dirty --always -long'. - The distance/hash is unconditional. - - Exceptions: - 1: no tags. HEX[-dirty] (note: no 'g' prefix) - """ - if pieces["closest-tag"]: - rendered = pieces["closest-tag"] - rendered += "-%d-g%s" % (pieces["distance"], pieces["short"]) - else: - # exception #1 - rendered = pieces["short"] - if pieces["dirty"]: - rendered += "-dirty" - return rendered - - -def render(pieces: Dict[str, Any], style: str) -> Dict[str, Any]: - """Render the given version pieces into the requested style.""" - if pieces["error"]: - return {"version": "unknown", - "full-revisionid": pieces.get("long"), - "dirty": None, - "error": pieces["error"], - "date": None} - - if not style or style == "default": - style = "pep440" # the default - - if style == "pep440": - rendered = render_pep440(pieces) - elif style == "pep440-branch": - rendered = render_pep440_branch(pieces) - elif style == "pep440-pre": - rendered = render_pep440_pre(pieces) - elif style == "pep440-post": - rendered = render_pep440_post(pieces) - elif style == "pep440-post-branch": - rendered = render_pep440_post_branch(pieces) - elif style == "pep440-old": - rendered = render_pep440_old(pieces) - elif style == "git-describe": - rendered = render_git_describe(pieces) - elif style == "git-describe-long": - rendered = render_git_describe_long(pieces) - else: - raise ValueError("unknown style '%s'" % style) - - return {"version": rendered, "full-revisionid": pieces["long"], - "dirty": pieces["dirty"], "error": None, - "date": pieces.get("date")} - - -def get_versions() -> Dict[str, Any]: - """Get version information or return default if unable to do so.""" - # I am in _version.py, which lives at ROOT/VERSIONFILE_SOURCE. If we have - # __file__, we can work backwards from there to the root. Some - # py2exe/bbfreeze/non-CPython implementations don't do __file__, in which - # case we can only use expanded keywords. - - cfg = get_config() - verbose = cfg.verbose - - try: - return git_versions_from_keywords(get_keywords(), cfg.tag_prefix, - verbose) - except NotThisMethod: - pass - - try: - root = os.path.realpath(__file__) - # versionfile_source is the relative path from the top of the source - # tree (where the .git directory might live) to this file. Invert - # this to find the root from __file__. - for _ in cfg.versionfile_source.split('/'): - root = os.path.dirname(root) - except NameError: - return {"version": "0+unknown", "full-revisionid": None, - "dirty": None, - "error": "unable to find root of source tree", - "date": None} - - try: - pieces = git_pieces_from_vcs(cfg.tag_prefix, root, verbose) - return render(pieces, cfg.style) - except NotThisMethod: - pass - - try: - if cfg.parentdir_prefix: - return versions_from_parentdir(cfg.parentdir_prefix, root, verbose) - except NotThisMethod: - pass - - return {"version": "0+unknown", "full-revisionid": None, - "dirty": None, - "error": "unable to compute version", "date": None} diff --git a/numba_dpex/core/parfors/compiler.py b/numba_dpex/core/parfors/compiler.py deleted file mode 100644 index 8864cc637c..0000000000 --- a/numba_dpex/core/parfors/compiler.py +++ /dev/null @@ -1,88 +0,0 @@ -# SPDX-FileCopyrightText: 2022 - 2024 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -from numba.core import compiler, ir -from numba.core import types as numba_types -from numba.core.compiler_lock import global_compiler_lock - -from numba_dpex.core import config -from numba_dpex.core.exceptions import ( - KernelHasReturnValueError, - UnreachableError, -) -from numba_dpex.core.pipelines.kernel_compiler import KernelCompiler - - -@global_compiler_lock -def compile_numba_ir_with_dpex( - pyfunc, - pyfunc_name, - args, - return_type, - target_context, - typing_context, - debug=False, - is_kernel=True, - extra_compile_flags=None, -): - """ - Compiles a function using class:`numba_dpex.core.pipelines.KernelCompiler` - and returns the compiled result. - - Args: - args: The list of arguments passed to the kernel. - debug (bool): Optional flag to turn on debug mode compilation. - extra_compile_flags: Extra flags passed to the compiler. - - Returns: - cres: Compiled result. - - Raises: - KernelHasReturnValueError: If the compiled function returns a - non-void value. - """ - # First compilation will trigger the initialization of the backend. - typingctx = typing_context - targetctx = target_context - - flags = compiler.Flags() - # Do not compile the function to a binary, just lower to LLVM - flags.debuginfo = config.DEBUGINFO_DEFAULT - flags.no_compile = True - flags.no_cpython_wrapper = True - flags.no_cfunc_wrapper = True - flags.nrt = False - - if debug: - flags.debuginfo = debug - - # Run compilation pipeline - if isinstance(pyfunc, ir.FunctionIR): - cres = compiler.compile_ir( - typingctx=typingctx, - targetctx=targetctx, - func_ir=pyfunc, - args=args, - return_type=return_type, - flags=flags, - locals={}, - pipeline_class=KernelCompiler, - ) - else: - raise UnreachableError() - - if ( - is_kernel - and cres.signature.return_type is not None - and cres.signature.return_type != numba_types.void - ): - raise KernelHasReturnValueError( - kernel_name=pyfunc_name, - return_type=cres.signature.return_type, - ) - # Linking depending libraries - library = cres.library - library.finalize() - - return cres diff --git a/numba_dpex/core/parfors/kernel_builder.py b/numba_dpex/core/parfors/kernel_builder.py index 6f059c8452..23aa46c7f9 100644 --- a/numba_dpex/core/parfors/kernel_builder.py +++ b/numba_dpex/core/parfors/kernel_builder.py @@ -3,25 +3,18 @@ # SPDX-License-Identifier: Apache-2.0 import copy -import sys import warnings -import dpctl -import dpctl.program as dpctl_prog from numba.core import ir, types from numba.core.errors import NumbaParallelSafetyWarning from numba.core.ir_utils import ( - add_offset_to_labels, get_name_var_table, get_unused_var_name, legalize_names, mk_unique_var, - remove_dead, remove_dels, - rename_labels, replace_var_names, ) -from numba.core.target_extension import target_override from numba.core.typing import signature from numba.parfors import parfor @@ -32,15 +25,11 @@ ) from numba_dpex.core.types.kernel_api.index_space_ids import ItemType from numba_dpex.core.utils.call_kernel_builder import SPIRVKernelModule -from numba_dpex.kernel_api_impl.spirv import spirv_generator from numba_dpex.kernel_api_impl.spirv.dispatcher import ( SPIRVKernelDispatcher, _SPIRVKernelCompileResult, ) -from ..descriptor import dpex_kernel_target -from ..types import DpnpNdArray -from .compiler import compile_numba_ir_with_dpex from .kernel_templates.range_kernel_template import RangeKernelTemplate diff --git a/numba_dpex/core/parfors/parfor_lowerer.py b/numba_dpex/core/parfors/parfor_lowerer.py index 3eb39b2cf5..4610cf9ef4 100644 --- a/numba_dpex/core/parfors/parfor_lowerer.py +++ b/numba_dpex/core/parfors/parfor_lowerer.py @@ -3,7 +3,6 @@ # SPDX-License-Identifier: Apache-2.0 import copy -from collections import namedtuple from llvmlite import ir as llvmir from numba.core import cgutils, ir, types diff --git a/numba_dpex/core/parfors/reduction_helper.py b/numba_dpex/core/parfors/reduction_helper.py index cb9e189977..eeeff3e486 100644 --- a/numba_dpex/core/parfors/reduction_helper.py +++ b/numba_dpex/core/parfors/reduction_helper.py @@ -17,10 +17,6 @@ from numba.parfors import parfor from numba.parfors.parfor_lowering_utils import ParforLoweringBuilder -from numba_dpex.core.datamodel.models import ( - dpex_data_model_manager as kernel_dmm, -) -from numba_dpex.core.utils.call_kernel_builder import KernelLaunchIRBuilder from numba_dpex.core.utils.cgutils_extra import get_llvm_type from numba_dpex.dpctl_iface import libsyclinterface_bindings as sycl diff --git a/numba_dpex/core/runtime/_dpexrt_python.c b/numba_dpex/core/runtime/_dpexrt_python.c index 10e441ed6c..5647a09241 100644 --- a/numba_dpex/core/runtime/_dpexrt_python.c +++ b/numba_dpex/core/runtime/_dpexrt_python.c @@ -45,12 +45,14 @@ static NRT_MemInfo *DPEXRT_MemInfo_fill(NRT_MemInfo *mi, bool value_is_float, int64_t value, const DPCTLSyclQueueRef qref); -static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj, +static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(NRT_api_functions *nrt, + PyObject *ndarrobj, void *data, npy_intp nitems, npy_intp itemsize, DPCTLSyclQueueRef qref); -static NRT_MemInfo *DPEXRT_MemInfo_alloc(npy_intp size, +static NRT_MemInfo *DPEXRT_MemInfo_alloc(NRT_api_functions *nrt, + npy_intp size, size_t usm_type, const DPCTLSyclQueueRef qref); static void usmndarray_meminfo_dtor(void *ptr, size_t size, void *info); @@ -58,7 +60,8 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct, int ndim, PyArray_Descr *descr); -static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj, +static int DPEXRT_sycl_usm_ndarray_from_python(NRT_api_functions *nrt, + PyObject *obj, usmarystruct_t *arystruct); static PyObject * DPEXRT_sycl_usm_ndarray_to_python_acqref(usmarystruct_t *arystruct, @@ -336,6 +339,11 @@ NRT_ExternalAllocator_new_for_usm(DPCTLSyclQueueRef qref, size_t usm_type) static void usmndarray_meminfo_dtor(void *ptr, size_t size, void *info) { MemInfoDtorInfo *mi_dtor_info = NULL; + // Warning: we are destructing sycl memory. MI destructor is called + // separately by numba. + DPEXRT_DEBUG(drt_debug_print("DPEXRT-DEBUG: Call to " + "usmndarray_meminfo_dtor at %s, line %d\n", + __FILE__, __LINE__)); // Sanity-check to make sure the mi_dtor_info is an actual pointer. if (!(mi_dtor_info = (MemInfoDtorInfo *)info)) { @@ -416,7 +424,8 @@ static MemInfoDtorInfo *MemInfoDtorInfo_new(NRT_MemInfo *mi, PyObject *owner) * of the dpnp.ndarray was allocated. * @return {return} A new NRT_MemInfo object */ -static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj, +static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(NRT_api_functions *nrt, + PyObject *ndarrobj, void *data, npy_intp nitems, npy_intp itemsize, @@ -427,8 +436,9 @@ static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj, MemInfoDtorInfo *midtor_info = NULL; DPCTLSyclContextRef cref = NULL; - // Allocate a new NRT_MemInfo object - if (!(mi = (NRT_MemInfo *)malloc(sizeof(NRT_MemInfo)))) { + // Allocate a new NRT_MemInfo object. By passing 0 we are just allocating + // MemInfo and not the `data` that the MemInfo object manages. + if (!(mi = (NRT_MemInfo *)nrt->allocate(0))) { DPEXRT_DEBUG(drt_debug_print( "DPEXRT-ERROR: Could not allocate a new NRT_MemInfo " "object at %s, line %d\n", @@ -505,7 +515,8 @@ static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj, * @return {return} A new NRT_MemInfo object, NULL if no NRT_MemInfo * object could be created. */ -static NRT_MemInfo *DPEXRT_MemInfo_alloc(npy_intp size, +static NRT_MemInfo *DPEXRT_MemInfo_alloc(NRT_api_functions *nrt, + npy_intp size, size_t usm_type, const DPCTLSyclQueueRef qref) { @@ -517,7 +528,7 @@ static NRT_MemInfo *DPEXRT_MemInfo_alloc(npy_intp size, "DPEXRT-DEBUG: Inside DPEXRT_MemInfo_alloc %s, line %d\n", __FILE__, __LINE__)); // Allocate a new NRT_MemInfo object - if (!(mi = (NRT_MemInfo *)malloc(sizeof(NRT_MemInfo)))) { + if (!(mi = (NRT_MemInfo *)nrt->allocate(0))) { DPEXRT_DEBUG(drt_debug_print( "DPEXRT-ERROR: Could not allocate a new NRT_MemInfo object.\n")); goto error; @@ -744,14 +755,18 @@ static struct PyUSMArrayObject *PyUSMNdArray_ARRAYOBJ(PyObject *obj) DPEXRT_DEBUG( drt_debug_print("DPEXRT-DEBUG: usm array was passed directly\n")); arrayobj = obj; + Py_INCREF(arrayobj); } else if (PyObject_HasAttrString(obj, "_array_obj")) { + // PyObject_GetAttrString gives reference arrayobj = PyObject_GetAttrString(obj, "_array_obj"); if (!arrayobj) return NULL; - if (!PyObject_TypeCheck(arrayobj, &PyUSMArrayType)) + if (!PyObject_TypeCheck(arrayobj, &PyUSMArrayType)) { + Py_DECREF(arrayobj); return NULL; + } } struct PyUSMArrayObject *pyusmarrayobj = @@ -791,7 +806,8 @@ static npy_intp product_of_shape(npy_intp *shape, npy_intp ndim) * instance of a dpnp.ndarray * @return {return} Error code representing success (0) or failure (-1). */ -static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj, +static int DPEXRT_sycl_usm_ndarray_from_python(NRT_api_functions *nrt, + PyObject *obj, usmarystruct_t *arystruct) { struct PyUSMArrayObject *arrayobj = NULL; @@ -803,17 +819,13 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj, PyGILState_STATE gstate; npy_intp itemsize = 0; - // Increment the ref count on obj to prevent CPython from garbage - // collecting the array. - // TODO: add extra description why do we need this - Py_IncRef(obj); - DPEXRT_DEBUG(drt_debug_print( "DPEXRT-DEBUG: In DPEXRT_sycl_usm_ndarray_from_python at %s, line %d\n", __FILE__, __LINE__)); // Check if the PyObject obj has an _array_obj attribute that is of // dpctl.tensor.usm_ndarray type. + // arrayobj is a new reference, reference of obj is borrowed if (!(arrayobj = PyUSMNdArray_ARRAYOBJ(obj))) { DPEXRT_DEBUG(drt_debug_print( "DPEXRT-ERROR: PyUSMNdArray_ARRAYOBJ check failed at %s, line %d\n", @@ -832,6 +844,7 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj, data = (void *)UsmNDArray_GetData(arrayobj); nitems = product_of_shape(shape, ndim); itemsize = (npy_intp)UsmNDArray_GetElementSize(arrayobj); + if (!(qref = UsmNDArray_GetQueueRef(arrayobj))) { DPEXRT_DEBUG(drt_debug_print( "DPEXRT-ERROR: UsmNDArray_GetQueueRef returned NULL at " @@ -841,7 +854,7 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj, } if (!(arystruct->meminfo = NRT_MemInfo_new_from_usmndarray( - obj, data, nitems, itemsize, qref))) + nrt, obj, data, nitems, itemsize, qref))) { DPEXRT_DEBUG(drt_debug_print( "DPEXRT-ERROR: NRT_MemInfo_new_from_usmndarray failed " @@ -850,6 +863,9 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj, goto error; } + Py_XDECREF(arrayobj); + Py_IncRef(obj); + arystruct->data = data; arystruct->sycl_queue = qref; arystruct->nitems = nitems; @@ -906,7 +922,7 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj, __FILE__, __LINE__)); gstate = PyGILState_Ensure(); // decref the python object - Py_DECREF(obj); + Py_XDECREF((PyObject *)arrayobj); // release the GIL PyGILState_Release(gstate); @@ -938,6 +954,7 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct, drt_debug_print("DPEXRT-DEBUG: In box_from_arystruct_parent.\n")); if (!(arrayobj = PyUSMNdArray_ARRAYOBJ(arystruct->parent))) { + Py_XDECREF(arrayobj); DPEXRT_DEBUG( drt_debug_print("DPEXRT-DEBUG: Arrayobj cannot be boxed from " "parent as parent pointer is NULL.\n")); @@ -945,6 +962,7 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct, } if ((void *)UsmNDArray_GetData(arrayobj) != arystruct->data) { + Py_XDECREF(arrayobj); DPEXRT_DEBUG(drt_debug_print( "DPEXRT-DEBUG: Arrayobj cannot be boxed " "from parent as data pointer in the arystruct is not the same as " @@ -952,12 +970,15 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct, return NULL; } - if (UsmNDArray_GetNDim(arrayobj) != ndim) + if (UsmNDArray_GetNDim(arrayobj) != ndim) { + Py_XDECREF(arrayobj); return NULL; + } p = arystruct->shape_and_strides; shape = UsmNDArray_GetShape(arrayobj); strides = UsmNDArray_GetStrides(arrayobj); + Py_XDECREF(arrayobj); // Ensure the shape of the array to be boxed matches the shape of the // original parent. diff --git a/numba_dpex/core/runtime/context.py b/numba_dpex/core/runtime/context.py index 80f0253101..4fe89abd7a 100644 --- a/numba_dpex/core/runtime/context.py +++ b/numba_dpex/core/runtime/context.py @@ -89,12 +89,14 @@ def meminfo_alloc_unchecked(self, builder, size, usm_type, queue_ref): mod = builder.module u64 = llvmir.IntType(64) fnty = llvmir.FunctionType( - cgutils.voidptr_t, [cgutils.intp_t, u64, cgutils.voidptr_t] + cgutils.voidptr_t, + [cgutils.voidptr_t, cgutils.intp_t, u64, cgutils.voidptr_t], ) fn = cgutils.get_or_insert_function(mod, fnty, "DPEXRT_MemInfo_alloc") fn.return_value.add_attribute("noalias") + nrt_api = self._context.nrt.get_nrt_api(builder) - ret = builder.call(fn, [size, usm_type, queue_ref]) + ret = builder.call(fn, [nrt_api, size, usm_type, queue_ref]) return ret @@ -168,13 +170,15 @@ def arraystruct_from_python(self, pyapi, obj, ptr): """ fnty = llvmir.FunctionType( - llvmir.IntType(32), [pyapi.pyobj, pyapi.voidptr] + llvmir.IntType(32), [pyapi.voidptr, pyapi.pyobj, pyapi.voidptr] ) + nrt_api = self._context.nrt.get_nrt_api(pyapi.builder) fn = pyapi._get_function(fnty, "DPEXRT_sycl_usm_ndarray_from_python") fn.args[0].add_attribute("nocapture") fn.args[1].add_attribute("nocapture") + fn.args[2].add_attribute("nocapture") - self.error = pyapi.builder.call(fn, (obj, ptr)) + self.error = pyapi.builder.call(fn, (nrt_api, obj, ptr)) return self.error diff --git a/numba_dpex/core/typing/dpnpdecl.py b/numba_dpex/core/typing/dpnpdecl.py index 5cf7fd78d9..720bada7d9 100644 --- a/numba_dpex/core/typing/dpnpdecl.py +++ b/numba_dpex/core/typing/dpnpdecl.py @@ -2,10 +2,7 @@ # # SPDX-License-Identifier: Apache-2.0 -import logging - import dpnp -import numpy as np from numba.core import types from numba.core.typing.npydecl import ( Numpy_rules_ufunc, @@ -25,7 +22,7 @@ def _install_operations(cls: NumpyRulesArrayOperator): for op, ufunc_name in cls._op_map.items(): infer_global(op)( type( - "NumpyRulesArrayOperator_" + ufunc_name, + "DpnpRulesArrayOperator_" + ufunc_name, (cls,), dict(key=op), ) @@ -35,36 +32,7 @@ def _install_operations(cls: NumpyRulesArrayOperator): class DpnpRulesArrayOperator(NumpyRulesArrayOperator): @property def ufunc(self): - try: - dpnpop = getattr(dpnp, self._op_map[self.key]) - npop = getattr(np, self._op_map[self.key]) - if not hasattr(dpnpop, "nin"): - dpnpop.nin = npop.nin - if not hasattr(dpnpop, "nout"): - dpnpop.nout = npop.nout - if not hasattr(dpnpop, "nargs"): - dpnpop.nargs = dpnpop.nin + dpnpop.nout - - # Check if the dpnp operation has a `types` attribute and if an - # AttributeError gets raised then "monkey patch" the attribute from - # numpy. If the attribute lookup raised a ValueError, it indicates - # that dpnp could not be resolve the supported types for the - # operation. Dpnp will fail to resolve the `types` if no SYCL - # devices are available on the system. For such a scenario, we print - # a user-level warning. - try: - dpnpop.types - except ValueError: - logging.exception( - f"The types attribute for the {dpnpop} fuction could not " - "be determined." - ) - except AttributeError: - dpnpop.types = npop.types - dpnpop.is_dpnp_ufunc = True - return dpnpop - except: - pass + return getattr(dpnp, self._op_map[self.key]) @classmethod def install_operations(cls): diff --git a/numba_dpex/dpnp_iface/dpnp_ufunc_db.py b/numba_dpex/dpnp_iface/dpnp_ufunc_db.py index 3cd1945595..f07bf47b03 100644 --- a/numba_dpex/dpnp_iface/dpnp_ufunc_db.py +++ b/numba_dpex/dpnp_iface/dpnp_ufunc_db.py @@ -3,7 +3,6 @@ # SPDX-License-Identifier: Apache-2.0 import copy -import logging import dpnp import numpy as np @@ -57,15 +56,9 @@ def _fill_ufunc_db_with_dpnp_ufuncs(ufunc_db): # variable is passed by value from numba.np.ufunc_db import _ufunc_db - failed_dpnpop_types_lst = [] for ufuncop in dpnpdecl.supported_ufuncs: if ufuncop == "erf": op = getattr(dpnp, "erf") - op.nin = 1 - op.nout = 1 - op.nargs = 2 - op.types = ["f->f", "d->d"] - op.is_dpnp_ufunc = True _unary_d_d = types.float64(types.float64) _unary_f_f = types.float32(types.float32) @@ -76,31 +69,7 @@ def _fill_ufunc_db_with_dpnp_ufuncs(ufunc_db): else: dpnpop = getattr(dpnp, ufuncop) npop = getattr(np, ufuncop) - if not hasattr(dpnpop, "nin"): - dpnpop.nin = npop.nin - if not hasattr(dpnpop, "nout"): - dpnpop.nout = npop.nout - if not hasattr(dpnpop, "nargs"): - dpnpop.nargs = dpnpop.nin + dpnpop.nout - - # Check if the dpnp operation has a `types` attribute and if an - # AttributeError gets raised then "monkey patch" the attribute from - # numpy. If the attribute lookup raised a ValueError, it indicates - # that dpnp could not be resolve the supported types for the - # operation. Dpnp will fail to resolve the `types` if no SYCL - # devices are available on the system. For such a scenario, we log - # dpnp operations for which the ValueError happened and print them - # as a user-level warning. It is done this way so that the failure - # to load the dpnpdecl registry due to the ValueError does not - # impede a user from importing numba-dpex. - try: - dpnpop.types - except ValueError: - failed_dpnpop_types_lst.append(ufuncop) - except AttributeError: - dpnpop.types = npop.types - - dpnpop.is_dpnp_ufunc = True + cp = copy.copy(_ufunc_db[npop]) ufunc_db.update({dpnpop: cp}) for key in list(ufunc_db[dpnpop].keys()): @@ -111,13 +80,3 @@ def _fill_ufunc_db_with_dpnp_ufuncs(ufunc_db): or "D->" in key ): ufunc_db[dpnpop].pop(key) - - if failed_dpnpop_types_lst: - try: - getattr(dpnp, failed_dpnpop_types_lst[0]).types - except ValueError: - ops = " ".join(failed_dpnpop_types_lst) - logging.exception( - "The types attribute for the following dpnp ops could not be " - f"determined: {ops}" - ) diff --git a/numba_dpex/kernel_api_impl/spirv/spirv_generator.py b/numba_dpex/kernel_api_impl/spirv/spirv_generator.py index c731a9c75d..192311f9e0 100644 --- a/numba_dpex/kernel_api_impl/spirv/spirv_generator.py +++ b/numba_dpex/kernel_api_impl/spirv/spirv_generator.py @@ -124,6 +124,7 @@ def finalize(self): "--spirv-ext=+SPV_EXT_shader_atomic_float_add", "--spirv-ext=+SPV_EXT_shader_atomic_float_min_max", "--spirv-ext=+SPV_INTEL_arbitrary_precision_integers", + "--spirv-ext=+SPV_INTEL_variable_length_array", ] for key in list(self.context.extra_compile_options.keys()): if key == LLVM_SPIRV_ARGS: diff --git a/numba_dpex/numba_patches/patch_is_ufunc.py b/numba_dpex/numba_patches/patch_is_ufunc.py deleted file mode 100644 index dee4a3f375..0000000000 --- a/numba_dpex/numba_patches/patch_is_ufunc.py +++ /dev/null @@ -1,23 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - - -def patch(): - """Patches the numba.np.ufunc.array_exprs._is_ufunc function to make it - possible to support dpnp universal functions (ufuncs). - - The extra condition is the check for the "is_dpnp_ufunc" attribute to - identify a non-NumPy ufunc. - """ - import numpy - from numba.np.ufunc.dufunc import DUFunc - - def _is_ufunc(func): - return isinstance(func, (numpy.ufunc, DUFunc)) or hasattr( - func, "is_dpnp_ufunc" - ) - - from numba.np.ufunc import array_exprs - - array_exprs._is_ufunc = _is_ufunc diff --git a/numba_dpex/numba_patches/patch_ufuncs.py b/numba_dpex/numba_patches/patch_ufuncs.py new file mode 100644 index 0000000000..e3e316dba2 --- /dev/null +++ b/numba_dpex/numba_patches/patch_ufuncs.py @@ -0,0 +1,94 @@ +# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 + +import logging + +import dpnp +import numpy as np + +from numba_dpex.core.typing import dpnpdecl + + +def patch(): + patch_is_ufunc() + patch_ufuncs() + + +def patch_is_ufunc(): + """Patches the numba.np.ufunc.array_exprs._is_ufunc function to make it + possible to support dpnp universal functions (ufuncs). + + The extra condition is the check for the "is_dpnp_ufunc" attribute to + identify a non-NumPy ufunc. + """ + import numpy + from numba.np.ufunc.dufunc import DUFunc + + def _is_ufunc(func): + return isinstance(func, (numpy.ufunc, DUFunc)) or hasattr( + func, "is_dpnp_ufunc" + ) + + from numba.np.ufunc import array_exprs + + array_exprs._is_ufunc = _is_ufunc + + +def patch_ufuncs(): + """Patches dpnp user functions to make them compatible with numpy, so we + can reuse numba's implementation. + + It adds "nin", "nout", "nargs" and "is_dpnp_ufunc" attributes to ufuncs. + """ + failed_dpnpop_types_lst = [] + + op = getattr(dpnp, "erf") + op.nin = 1 + op.nout = 1 + op.nargs = 2 + op.types = ["f->f", "d->d"] + op.is_dpnp_ufunc = True + + for ufuncop in dpnpdecl.supported_ufuncs: + if ufuncop == "erf": + continue + + dpnpop = getattr(dpnp, ufuncop) + npop = getattr(np, ufuncop) + + if not hasattr(dpnpop, "nin"): + dpnpop.nin = npop.nin + if not hasattr(dpnpop, "nout"): + dpnpop.nout = npop.nout + if not hasattr(dpnpop, "nargs"): + dpnpop.nargs = dpnpop.nin + dpnpop.nout + + # Check if the dpnp operation has a `types` attribute and if an + # AttributeError gets raised then "monkey patch" the attribute from + # numpy. If the attribute lookup raised a ValueError, it indicates + # that dpnp could not be resolve the supported types for the + # operation. Dpnp will fail to resolve the `types` if no SYCL + # devices are available on the system. For such a scenario, we log + # dpnp operations for which the ValueError happened and print them + # as a user-level warning. It is done this way so that the failure + # to load the dpnpdecl registry due to the ValueError does not + # impede a user from importing numba-dpex. + try: + dpnpop.types + except ValueError: + failed_dpnpop_types_lst.append(ufuncop) + except AttributeError: + dpnpop.types = npop.types + + dpnpop.is_dpnp_ufunc = True + + if len(failed_dpnpop_types_lst) > 0: + try: + getattr(dpnp, failed_dpnpop_types_lst[0]).types + except ValueError: + ops = " ".join(failed_dpnpop_types_lst) + logging.exception( + "The types attribute for the following dpnp ops could not be " + f"determined: {ops}" + ) diff --git a/numba_dpex/tests/_helper.py b/numba_dpex/tests/_helper.py index c790cb5b73..41b31518be 100644 --- a/numba_dpex/tests/_helper.py +++ b/numba_dpex/tests/_helper.py @@ -122,7 +122,7 @@ def is_windows(): ] skip_no_gdb = pytest.mark.skipif( - config.TESTING_SKIP_NO_DEBUGGING and not shutil.which("gdb-oneapi"), + config.TESTING_SKIP_NO_DEBUGGING or not shutil.which("gdb-oneapi"), reason="Intel® Distribution for GDB* is not available", ) diff --git a/numba_dpex/tests/codegen/test_intenum_literal_codegen.py b/numba_dpex/tests/codegen/test_intenum_literal_codegen.py index 5d875f37af..ef19eb2bed 100644 --- a/numba_dpex/tests/codegen/test_intenum_literal_codegen.py +++ b/numba_dpex/tests/codegen/test_intenum_literal_codegen.py @@ -47,7 +47,7 @@ def pass_flags_to_func(a): pattern = re.compile( r"call spir_func i32 @\_Z.*bitwise\_or" - r"\_flags.*\(i64\* nonnull %.*, i64 1, i64 2\)" + r"\_flags.*\(i64\*\s(\w+)?\s*%.*, i64 1, i64 2\)" ) assert re.search(pattern, llvm_ir_mod) is not None diff --git a/numba_dpex/tests/debugging/gdb.py b/numba_dpex/tests/debugging/gdb.py index f47fad117b..b743507c15 100644 --- a/numba_dpex/tests/debugging/gdb.py +++ b/numba_dpex/tests/debugging/gdb.py @@ -33,8 +33,8 @@ r"layout=[A-Z],\s+" r"address_space=[0-4],\s+" r"usm_type=[a-z]+,\s+" - r"device=[a-z:0-9]+,\s+" - r"sycl_queue=[A-Za-z:0-9 ]+\)" + r"device=[a-z\_:0-9]+,\s+" + r"sycl_queue=[A-Za-z:0-9\s\_:]+\)" ) @@ -52,7 +52,7 @@ def spawn(self): env["NUMBA_DEBUGINFO"] = "1" self.child = pexpect.spawn( - "gdb-oneapi -q python", env=env, encoding="utf-8" + "gdb-oneapi -q python", env=env, encoding="utf-8", timeout=60 ) if config.TESTING_LOG_DEBUGGING: self.child.logfile = sys.stdout @@ -109,7 +109,7 @@ def expect_eol(self): self.child.expect(r"[^\n]*\n") def expect_hit_breakpoint(self, expected_location=None): - expect = r"Thread [0-9A-Za-z \"]+ hit Breakpoint [0-9\.]+" + expect = r"Breakpoint [0-9\.]+" if expected_location is not None: # function name + args could be long, so we have to assume that # the message may be splitted in multiple lines. It potentially can diff --git a/numba_dpex/tests/debugging/test_breakpoints.py b/numba_dpex/tests/debugging/test_breakpoints.py index 701d4b823b..f76411b100 100644 --- a/numba_dpex/tests/debugging/test_breakpoints.py +++ b/numba_dpex/tests/debugging/test_breakpoints.py @@ -58,7 +58,7 @@ def test_device_func_breakpoint( app.breakpoint(breakpoint, condition=condition) app.run(f"side-by-side.py --api={api}") - app.expect_hit_breakpoint("side-by-side.py:15") + app.expect_hit_breakpoint(expected_location="side-by-side.py:15") if exp_var is not None: app.print(exp_var, expected=exp_val) diff --git a/numba_dpex/tests/debugging/test_info.py b/numba_dpex/tests/debugging/test_info.py index c9f75dd102..693f487fcb 100644 --- a/numba_dpex/tests/debugging/test_info.py +++ b/numba_dpex/tests/debugging/test_info.py @@ -78,7 +78,7 @@ def test_info_args( ): app.breakpoint(breakpoint) app.run(script) - app.expect_hit_breakpoint(breakpoint) + app.expect_hit_breakpoint(expected_location=breakpoint) app.expect(expected_line, with_eol=True) if kind == "info": @@ -99,7 +99,7 @@ def test_info_args( def test_info_functions(app): app.breakpoint("simple_sum.py:12") app.run("simple_sum.py") - app.expect_hit_breakpoint("simple_sum.py:12") + app.expect_hit_breakpoint(expected_location="simple_sum.py:12") app.expect(r"12\s+i = item.get_id\(0\)", with_eol=True) app.info_functions("data_parallel_sum") @@ -119,7 +119,7 @@ def test_print_array_element(app, api): app.breakpoint("side-by-side-2.py:17", condition="param_a == 5") app.run(f"side-by-side-2.py --api={api}") - app.expect_hit_breakpoint("side-by-side-2.py:17") + app.expect_hit_breakpoint(expected_location="side-by-side-2.py:17") # We can access only c_array, not python style array app.print("b.data[5]", 5) @@ -142,7 +142,7 @@ def test_print_array_element(app, api): def test_assignment_to_variable(app, api, assign): app.breakpoint("side-by-side-2.py:17", condition="param_a == 5") app.run(f"side-by-side-2.py --api={api}") - app.expect_hit_breakpoint("side-by-side-2.py:17") + app.expect_hit_breakpoint(expected_location="side-by-side-2.py:17") app.print("param_a", expected=5) if assign == "print": diff --git a/numba_dpex/tests/debugging/test_stepping.py b/numba_dpex/tests/debugging/test_stepping.py index 71fedcab83..3e802a1db0 100644 --- a/numba_dpex/tests/debugging/test_stepping.py +++ b/numba_dpex/tests/debugging/test_stepping.py @@ -21,7 +21,7 @@ def test_next(app: gdb): app.breakpoint("simple_dpex_func.py:18") app.run("simple_dpex_func.py") - app.expect_hit_breakpoint("simple_dpex_func.py:18") + app.expect_hit_breakpoint(expected_location="simple_dpex_func.py:18") app.expect(r"18\s+i = item.get_id\(0\)", with_eol=True) app.set_scheduler_lock() app.next() @@ -37,7 +37,7 @@ def test_next(app: gdb): def test_step(app: gdb): app.breakpoint("simple_dpex_func.py:19") app.run("simple_dpex_func.py") - app.expect_hit_breakpoint("simple_dpex_func.py:19") + app.expect_hit_breakpoint(expected_location="simple_dpex_func.py:19") app.expect( r"19\s+c_in_kernel\[i\] = func_sum\(a_in_kernel\[i\], b_in_kernel\[i\]\)", with_eol=True, @@ -57,7 +57,7 @@ def test_step(app: gdb): def test_stepi(app: gdb, func: str): app.breakpoint("simple_dpex_func.py:19") app.run("simple_dpex_func.py") - app.expect_hit_breakpoint("simple_dpex_func.py:19") + app.expect_hit_breakpoint(expected_location="simple_dpex_func.py:19") app.expect( r"19\s+c_in_kernel\[i\] = func_sum\(a_in_kernel\[i\], b_in_kernel\[i\]\)", with_eol=True, diff --git a/pyproject.toml b/pyproject.toml index e44e25b640..aa7cc94eb3 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -1,15 +1,21 @@ [build-system] build-backend = "setuptools.build_meta" requires = [ - # TODO: keep in sync with [project.dependencies] and /conda-recipe/meta.yaml + # TODO: keep in sync with [project.dependencies] + "wheel>=0.43", + "build>=1.1", "setuptools>=63.0.0", "scikit-build>=0.17.0", "ninja>=1.11.1; platform_system!='Windows'", "cmake>=3.29.0", - "numba>=0.59.0", + # We need dpctl for UsmNdArray integration for dpcpp code "dpctl>=0.16.1", - "numpy>=1.24.0", - "wheel", + # We need numba for runtime cpp headers + "numba>=0.59.0,<0.60.0a0", + "llvmlite>=0.42.0", + # Do we need dpnp at build time? + "dpnp >=0.14", + "numpy >=1.24", # WARNING: check with doc how to upgrade "versioneer[toml]==0.29" ] @@ -32,14 +38,16 @@ classifiers = [ dependencies = [ # TODO: keep in sync with [build-system.requires] and /conda-recipe/meta.yaml # This restrictions are for dependabot, actual restrictions are set with - # conda. - # TODO: populate it during build process + # conda. TODO: populate it during build process + # TODO: do we have to set sycl runtime dependencies here + # "dpcpp-cpp-rt>=0.59.0", + # "intel-cmplr-lib-rt>=0.59.0" "numba>=0.59.0", "llvmlite>=0.42.0", "dpctl>=0.16.1", "dpnp>=0.14.0", "numpy>=1.24.0", - "dpcpp_llvm_spirv>=2024.0" + "dpcpp_llvm_spirv>=2024.1" ] description = "An extension for Numba to add data-parallel offload capability" dynamic = ["version"] diff --git a/setup.py b/setup.py index a5613c1884..c545d3d1a4 100644 --- a/setup.py +++ b/setup.py @@ -3,10 +3,11 @@ # SPDX-License-Identifier: Apache-2.0 +import os import re import versioneer -from setuptools import find_packages +from setuptools import find_namespace_packages, find_packages from skbuild import setup """Top level setup.py file. Uses scikit-build. @@ -50,6 +51,16 @@ def to_cmake_format(version: str): # Get the project version __version__ = versioneer.get_version() +# Keep same style as versioneer in case ve would want to switch back +with open(os.path.join("numba_dpex", "_version.py"), "w") as fh: + print( + "# SPDX-FileCopyrightText: 2024 Intel Corporation\n#\n" + "# SPDX-License-Identifier: Apache-2.0\n\n\ndef get_versions():\n" + ' return {"version": "' + __version__ + '"}', + file=fh, + ) + + # Main setup setup( version=__version__, @@ -58,7 +69,8 @@ def to_cmake_format(version: str): # Must be passed vis setup.py: # https://github.com/scikit-build/scikit-build/issues/864 # TODO: switch to pyproject toml after switching to scikit-build-core - packages=find_packages("."), + packages=find_packages(".") + + find_namespace_packages(".", include=["numba_dpex.examples.*"]), # Needs for examples. # TODO: change to false once move examples out of package. include_package_data=True,