diff --git a/.github/workflows/bump-sycl-deps.yml b/.github/workflows/bump-sycl-deps.yml new file mode 100644 index 000000000000..582caf709905 --- /dev/null +++ b/.github/workflows/bump-sycl-deps.yml @@ -0,0 +1,139 @@ +name: Bump open source SYCL build deps + +# Keeps the pinned oclcpuexp driver (intel/llvm YYYY-WWNN release) and oneTBB +# release used by os-llvm-sycl-build.yml up to date by opening a pull request +# when a newer upstream release is available. + +on: + # Check weekly on Saturday, and allow manual runs. + schedule: + - cron: "0 6 * * 6" + workflow_dispatch: + +permissions: read-all + +# Only one bump run at a time; a newer run supersedes an in-progress one. +concurrency: + group: ${{ github.workflow }}-${{ github.ref }} + cancel-in-progress: true + +jobs: + bump: + name: Bump oclcpuexp and oneTBB pins + + runs-on: ubuntu-latest + timeout-minutes: 15 + + permissions: + contents: write + pull-requests: write + + env: + WF: .github/workflows/os-llvm-sycl-build.yml + GH_TOKEN: ${{ github.token }} + + steps: + - name: Checkout repo + uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 + + - name: Find latest intel/llvm driver release tag + id: driver + uses: oprypin/find-latest-tag@6957ac556fa6d349727ecabfcaaf9f8e5ee37124 # v1.1.3 + with: + repository: intel/llvm + releases-only: true + # Driver releases are tagged like "2026-WW28"; ignore nightly-* tags. + regex: '^\d{4}-WW\d+$' + # Pick the most recently published matching release. Must stay false: + # sort-tags: true would semver-sort and mis-parse the YYYY-WWNN tag. + sort-tags: false + + - name: Find latest oneTBB release tag + id: tbb + uses: oprypin/find-latest-tag@6957ac556fa6d349727ecabfcaaf9f8e5ee37124 # v1.1.3 + with: + repository: uxlfoundation/oneTBB + releases-only: true + regex: '^v\d+\.\d+\.\d+$' + # Most recently published matching release (see note above). + sort-tags: false + + - name: Resolve asset names and derive values + id: resolve + env: + DRIVER_TAG: ${{ steps.driver.outputs.tag }} + TBB_TAG: ${{ steps.tbb.outputs.tag }} + run: | + set -euo pipefail + + # The oclcpuexp filename embeds its own version and is not derivable + # from the tag, so read it back from the driver release assets. + OCLCPUEXP_FN=$(gh api "repos/intel/llvm/releases/tags/${DRIVER_TAG}" \ + --jq '.assets[].name' \ + | grep -E '^oclcpuexp-.*_rel\.tar\.gz$' | head -1) + if [[ -z "${OCLCPUEXP_FN}" ]]; then + echo "::error::no oclcpuexp asset found in intel/llvm release ${DRIVER_TAG}" + exit 1 + fi + + # oneTBB filenames are derived from the tag (e.g. v2023.1.0). + TBB_VER=${TBB_TAG#v} + + { + echo "driver_tag=${DRIVER_TAG}" + echo "oclcpuexp_fn=${OCLCPUEXP_FN}" + echo "tbb_tag=${TBB_TAG}" + echo "tbb_fn=oneapi-tbb-${TBB_VER}-lin.tgz" + echo "tbb_dir=oneapi-tbb-${TBB_VER}" + } >> "${GITHUB_OUTPUT}" + + - name: Apply pins to workflow file + env: + DRIVER_TAG: ${{ steps.resolve.outputs.driver_tag }} + OCLCPUEXP_FN: ${{ steps.resolve.outputs.oclcpuexp_fn }} + TBB_TAG: ${{ steps.resolve.outputs.tbb_tag }} + TBB_FN: ${{ steps.resolve.outputs.tbb_fn }} + TBB_DIR: ${{ steps.resolve.outputs.tbb_dir }} + run: | + set -euo pipefail + # Surgical, comment-preserving edits: rewrite only the value after each + # key and keep indentation. "|" delimiter so the "/"-heavy URL is safe. + sed -i -E \ + -e "s|^( *DRIVER_PATH: ).*|\1${DRIVER_TAG}|" \ + -e "s|^( *OCLCPUEXP_FN: ).*|\1${OCLCPUEXP_FN}|" \ + -e "s|^( *TBB_URL: ).*|\1https://github.com/uxlfoundation/oneTBB/releases/download/${TBB_TAG}/|" \ + -e "s|^( *TBB_INSTALL_DIR: ).*|\1${TBB_DIR}|" \ + -e "s|^( *TBB_FN: ).*|\1${TBB_FN}|" \ + "${WF}" + echo "---- resulting diff ----" + git diff -- "${WF}" || true + + - name: Create pull request + uses: peter-evans/create-pull-request@5f6978faf089d4d20b00c7766989d076bb2fc7f1 # v8.1.1 + with: + token: ${{ github.token }} + add-paths: ${{ env.WF }} + branch: bot/bump-sycl-build-deps + delete-branch: true + labels: autoupdate + commit-message: | + Bump oclcpuexp/oneTBB pins in os-llvm-sycl-build workflow + + driver: ${{ steps.resolve.outputs.driver_tag }} / ${{ steps.resolve.outputs.oclcpuexp_fn }} + oneTBB: ${{ steps.resolve.outputs.tbb_tag }} + title: "Bump open source SYCL build dependencies" + body: | + Automated bump of the pinned open source SYCL build dependencies in + `os-llvm-sycl-build.yml`. + + | dependency | new value | + | --- | --- | + | intel/llvm driver release | `${{ steps.resolve.outputs.driver_tag }}` | + | oclcpuexp | `${{ steps.resolve.outputs.oclcpuexp_fn }}` | + | oneTBB | `${{ steps.resolve.outputs.tbb_tag }}` | + + The nightly DPC++ compiler is discovered at run time and is not + affected by this change. + + Trigger the **Build with Open Source LLVM SYCL compiler** workflow on + this branch to validate before merging. diff --git a/.github/workflows/os-llvm-sycl-build.yml b/.github/workflows/os-llvm-sycl-build.yml new file mode 100644 index 000000000000..8c8ecdf935fe --- /dev/null +++ b/.github/workflows/os-llvm-sycl-build.yml @@ -0,0 +1,201 @@ +name: Build with Open Source LLVM SYCL compiler + +on: + push: + branches: + - master + pull_request: + # To be able to be triggered manually + workflow_dispatch: + +permissions: read-all + +# Only one build at a time per ref; a newer run supersedes an in-progress one. +concurrency: + group: ${{ github.workflow }}-${{ github.ref }} + cancel-in-progress: true + +jobs: + install-compiler: + name: Build with nightly build of DPC++ toolchain + + runs-on: ubuntu-latest + timeout-minutes: 120 + + defaults: + run: + shell: bash -el {0} + + env: + DOWNLOAD_URL_PREFIX: https://github.com/intel/llvm/releases/download + DRIVER_PATH: 2026-WW28 + OCLCPUEXP_FN: oclcpuexp-2026.22.6.1.17_160000_rel.tar.gz + TBB_URL: https://github.com/uxlfoundation/oneTBB/releases/download/v2023.1.0/ + TBB_INSTALL_DIR: oneapi-tbb-2023.1.0 + TBB_FN: oneapi-tbb-2023.1.0-lin.tgz + dpctl-repo-path: '${{ github.workspace }}/dpctl' + + steps: + - name: Install hwloc + run: sudo apt install hwloc + + - name: Install lld + run: sudo apt install lld + + - name: Download and install nightly and components + env: + ARTIFACT_NAME: sycl_linux + USE_LATEST_SYCLOS: 1 + 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 nightly tags and reverse order + LLVM_TAGS=$(git -c 'versionsort.suffix=-' ls-remote --tags --sort='v:refname' https://github.com/intel/llvm.git | \ + grep 'refs/tags/nightly-' | awk '{a[i++]=$0} END {for (j=i-1; j>=0;) print a[j--] }') + export LLVM_TAGS + # initialize + unset DEPLOY_NIGHTLY_TAG + unset DEPLOY_LLVM_TAG_SHA + + # go through tags and find the most recent one where nighly build binary is available + while IFS= read -r NEXT_LLVM_TAG; do + NEXT_LLVM_TAG_SHA=$(echo "${NEXT_LLVM_TAG}" | awk '{print $1}') + export NEXT_LLVM_TAG_SHA + 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}')") + export NEXT_NIGHTLY_TAG + if wget -S --spider "${DOWNLOAD_URL_PREFIX}/${NEXT_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz" 2>&1 | grep -q 'HTTP/1.1 200 OK'; + then + DEPLOY_NIGHTLY_TAG="${NEXT_NIGHTLY_TAG}" + export DEPLOY_NIGHTLY_TAG + DEPLOY_LLVM_TAG_SHA="${NEXT_LLVM_TAG_SHA}" + export DEPLOY_LLVM_TAG_SHA + break + fi + done <<< "${LLVM_TAGS}" + else + # Use latest known to work tag instead + DEPLOY_NIGHTLY_TAG="sycl-nightly%2F20230606" + export DEPLOY_NIGHTLY_TAG + DEPLOY_LLVM_TAG_SHA=f44d0133d4b0077298f034697a1f3818ff1d6134 + export DEPLOY_LLVM_TAG_SHA + fi + + [[ -n "${DEPLOY_NIGHTLY_TAG}" ]] || exit 1 + [[ -n "${DEPLOY_LLVM_TAG_SHA}" ]] || exit 1 + echo "Using ${DEPLOY_NIGHTLY_TAG} corresponding to intel/llvm at ${DEPLOY_LLVM_TAG_SHA}" + + rm -rf "${ARTIFACT_NAME}.tar.gz" + wget "${DOWNLOAD_URL_PREFIX}/${DEPLOY_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz" + wget "${DOWNLOAD_URL_PREFIX}/${DRIVER_PATH}/${OCLCPUEXP_FN}" + wget "${TBB_URL}/${TBB_FN}" + mkdir -p dpcpp_compiler + tar xf "${ARTIFACT_NAME}.tar.gz" -C dpcpp_compiler + mkdir -p oclcpuexp + tar xf "${OCLCPUEXP_FN}" -C oclcpuexp + tar xf "${TBB_FN}" + cp oclcpuexp/x64/libOpenCL.so* dpcpp_compiler/lib/ + + - name: Setup Python + uses: actions/setup-python@5fda3b95a4ea91299a34e894583c3862153e4b97 # v7.0.0 + with: + python-version: '3.14' + + - name: Checkout repo + uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 + with: + fetch-depth: 0 + + - name: Checkout dpctl repo + uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 + with: + repository: IntelPython/dpctl + fetch-depth: 0 + path: ${{ env.dpctl-repo-path }} + + - name: Install dpnp dependencies + run: | + pip install numpy cython setuptools"<80" pytest scikit-build cmake ninja versioneer[toml]==0.29 \ + mkl-devel-dpcpp onedpl-devel + + - name: Create set_allvars.sh + run: | + cat << 'EOF' > set_allvars.sh + #!/usr/bin/bash + export SYCL_BUNDLE_FOLDER=/home/runner/work/sycl_bundle + 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}/${TBB_INSTALL_DIR}/lib/intel64/gcc4.8:${LD_LIBRARY_PATH} + export OCL_ICD_VENDORS= + export OCL_ICD_FILENAMES=libintelocl.so + # Help CMake find MKL/oneDPL/TBB shipped in the Python environment. + PYTHON_PREFIX=$(python -c "import sys, os; print(os.path.dirname(os.path.dirname(sys.executable)))") + export MKLROOT=${PYTHON_PREFIX} + export oneDPL_ROOT=${PYTHON_PREFIX} + export TBBROOT=${SYCL_BUNDLE_FOLDER}/${TBB_INSTALL_DIR} + # The open source DPC++ compiler does not bundle oneDPL/MKL headers. + # dpnp extensions include e.g. without linking the + # oneDPL target, so expose the pip package headers on the default + # compiler include path. Use the *_INCLUDE_PATH vars (treated as + # -isystem) rather than CPATH (-I) so deprecation warnings from these + # headers stay suppressed. + export C_INCLUDE_PATH=${PYTHON_PREFIX}/include:${C_INCLUDE_PATH} + export CPLUS_INCLUDE_PATH=${PYTHON_PREFIX}/include:${CPLUS_INCLUDE_PATH} + EOF + chmod +x set_allvars.sh + cat set_allvars.sh + + - name: Report compiler version + run: | + source set_allvars.sh + clang++ --version + + - name: Run sycl-ls + run: | + source set_allvars.sh + sycl-ls + + - name: Build and install dpctl + run: | + source "${GITHUB_WORKSPACE}/set_allvars.sh" + # Build dpctl from source with the same nightly so both + # share one SYCL runtime. + # + # Install non-editable so scikit-build runs CMake's install step and + # produces a proper package layout (include/, resources/cmake with the + # config) that dpnp's find_package(Dpctl) needs. --no-deps avoids + # pulling intel-sycl-rt from PyPI, which would reintroduce the + # conflicting released libsycl (the nightly one is on LD_LIBRARY_PATH). + CC=clang CXX=clang++ pip install --no-build-isolation --no-deps . || exit 1 + working-directory: ${{ env.dpctl-repo-path }} + + - name: Build dpnp + run: | + source set_allvars.sh + # Build against the oneMath interface, so the math layer is compiled + # with the same open source nightly compiler. + # + # The open source DPC++ nightly ships an IntelSYCLConfig.cmake that + # lacks add_sycl_to_target(), so force the vendored module which + # provides it. + python scripts/build_locally.py --c-compiler=clang --cxx-compiler=clang++ \ + --compiler-root="${SYCL_BUNDLE_FOLDER}/dpcpp_compiler/bin" \ + --onemath \ + --cmake-opts="-DIntelSYCL_DIR=${GITHUB_WORKSPACE}/dpnp/backend/cmake/Modules" || exit 1 + + - name: Smoke test + run: | + source set_allvars.sh + python -m dpctl -f + python -c "import dpnp; print(dpnp.__version__)" + + - name: Run dpnp/tests + env: + SYCL_CACHE_PERSISTENT: 1 + SKIP_TENSOR_TESTS: 0 + run: | + source set_allvars.sh + python -m pytest -ra dpnp/tests diff --git a/dpnp/backend/cmake/Modules/IntelSYCLConfig.cmake b/dpnp/backend/cmake/Modules/IntelSYCLConfig.cmake index 7ba269c70c08..d31816871c1f 100644 --- a/dpnp/backend/cmake/Modules/IntelSYCLConfig.cmake +++ b/dpnp/backend/cmake/Modules/IntelSYCLConfig.cmake @@ -172,6 +172,18 @@ function(parse_compiler_version compiler_name version_number) # Parse Intel Compiler Version string(REGEX REPLACE "Intel\\(R\\) (.*) Compiler ([0-9]+\\.[0-9]+\\.[0-9]+) (.*)" "\\2" SYCL_VERSION_STRING_MATCH ${INTEL_VERSION_STRING}) + else() + # Open source DPC++/LLVM SYCL compiler reports e.g. + # "clang version 23.0.0git (https://github.com/intel/llvm ...)" + string(REGEX MATCH "clang version ([0-9]+\\.[0-9]+\\.[0-9]+)" + CLANG_VERSION_STRING ${COMPILER_VERSION_STRING}) + if(CLANG_VERSION_STRING) + string(REGEX REPLACE "clang version ([0-9]+\\.[0-9]+\\.[0-9]+)" "\\1" + SYCL_VERSION_STRING_MATCH ${CLANG_VERSION_STRING}) + endif() + endif() + + if(SYCL_VERSION_STRING_MATCH) string(REPLACE "." ";" SYCL_VERSION_LIST ${SYCL_VERSION_STRING_MATCH}) list(GET SYCL_VERSION_LIST 0 VERSION_MAJOR) list(GET SYCL_VERSION_LIST 1 VERSION_MINOR) diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp index de28af11c2b7..cc54cf373aa4 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp @@ -87,7 +87,7 @@ struct OutputType static int output_typeid_vector[td_ns::num_types]; template -struct TypeMapFactory +struct ErfTypeMapFactory { std::enable_if_t::value, int> get() { @@ -223,7 +223,7 @@ MACRO_DEFINE_IMPL(erfcinv, Erfcinv); void init_erf_funcs(py::module_ m) { using impl::output_typeid_vector; - init_dispatch_vector(output_typeid_vector); + init_dispatch_vector(output_typeid_vector); auto erf_result_type_pyapi = [&](const py::dtype &dtype) { return py_int::py_unary_ufunc_result_type(dtype, output_typeid_vector); diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp index 452bd1a71d4c..8967069e73da 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp @@ -85,7 +85,7 @@ struct OutputType static int float_power_output_typeid_table[td_ns::num_types][td_ns::num_types]; template -struct TypeMapFactory +struct FloatPowerTypeMapFactory { std::enable_if_t::value, int> get() { @@ -96,7 +96,8 @@ struct TypeMapFactory static void populate_float_power_dispatch_tables(void) { - init_dispatch_table(float_power_output_typeid_table); + init_dispatch_table( + float_power_output_typeid_table); } } // namespace impl diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp index 2971c3eb4aca..d8cfb8cf75ba 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp @@ -59,7 +59,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct ContigFactory \ + struct __name__##_ContigFactory \ { \ fnT get() \ { \ @@ -76,7 +76,7 @@ namespace ext_ns = ext::common; }; \ \ template \ - struct TypeMapFactory \ + struct __name__##_TypeMapFactory \ { \ std::enable_if_t::value, int> get() \ { \ @@ -103,7 +103,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct StridedFactory \ + struct __name__##_StridedFactory \ { \ fnT get() \ { \ @@ -122,12 +122,12 @@ namespace ext_ns = ext::common; void populate_##__name__##_dispatch_vectors(void) \ { \ ext_ns::init_dispatch_vector( \ + __name__##_ContigFactory>( \ __name__##_contig_dispatch_vector); \ ext_ns::init_dispatch_vector( \ + __name__##_StridedFactory>( \ __name__##_strided_dispatch_vector); \ - ext_ns::init_dispatch_vector( \ + ext_ns::init_dispatch_vector( \ __name__##_output_typeid_vector); \ }; @@ -151,7 +151,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct ContigFactory \ + struct __name__##_ContigFactory \ { \ fnT get() \ { \ @@ -170,7 +170,7 @@ namespace ext_ns = ext::common; }; \ \ template \ - struct TypeMapFactory \ + struct __name__##_TypeMapFactory \ { \ std::enable_if_t>::value, \ std::pair> \ @@ -202,7 +202,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct StridedFactory \ + struct __name__##_StridedFactory \ { \ fnT get() \ { \ @@ -223,12 +223,13 @@ namespace ext_ns = ext::common; void populate_##__name__##_dispatch_vectors(void) \ { \ ext_ns::init_dispatch_vector( \ + __name__##_ContigFactory>( \ __name__##_contig_dispatch_vector); \ ext_ns::init_dispatch_vector( \ + __name__##_StridedFactory>( \ __name__##_strided_dispatch_vector); \ - ext_ns::init_dispatch_vector, TypeMapFactory>( \ + ext_ns::init_dispatch_vector, \ + __name__##_TypeMapFactory>( \ __name__##_output_typeid_vector); \ }; @@ -256,7 +257,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct ContigFactory \ + struct __name__##_ContigFactory \ { \ fnT get() \ { \ @@ -275,7 +276,7 @@ namespace ext_ns = ext::common; }; \ \ template \ - struct TypeMapFactory \ + struct __name__##_TypeMapFactory \ { \ std::enable_if_t::value, int> get() \ { \ @@ -305,7 +306,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct StridedFactory \ + struct __name__##_StridedFactory \ { \ fnT get() \ { \ @@ -325,12 +326,12 @@ namespace ext_ns = ext::common; void populate_##__name__##_dispatch_tables(void) \ { \ ext_ns::init_dispatch_table( \ + __name__##_ContigFactory>( \ __name__##_contig_dispatch_table); \ ext_ns::init_dispatch_table( \ + __name__##_StridedFactory>( \ __name__##_strided_dispatch_table); \ - ext_ns::init_dispatch_table( \ + ext_ns::init_dispatch_table( \ __name__##_output_typeid_table); \ }; @@ -358,7 +359,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct ContigFactory \ + struct __name__##_ContigFactory \ { \ fnT get() \ { \ @@ -380,7 +381,7 @@ namespace ext_ns = ext::common; }; \ \ template \ - struct TypeMapFactory \ + struct __name__##_TypeMapFactory \ { \ std::enable_if_t>::value, \ std::pair> \ @@ -415,7 +416,7 @@ namespace ext_ns = ext::common; } \ \ template \ - struct StridedFactory \ + struct __name__##_StridedFactory \ { \ fnT get() \ { \ @@ -438,11 +439,12 @@ namespace ext_ns = ext::common; void populate_##__name__##_dispatch_tables(void) \ { \ ext_ns::init_dispatch_table( \ + __name__##_ContigFactory>( \ __name__##_contig_dispatch_table); \ ext_ns::init_dispatch_table( \ + __name__##_StridedFactory>( \ __name__##_strided_dispatch_table); \ - ext_ns::init_dispatch_table, TypeMapFactory>( \ + ext_ns::init_dispatch_table, \ + __name__##_TypeMapFactory>( \ __name__##_output_typeid_table); \ }; diff --git a/dpnp/backend/extensions/vm/common.hpp b/dpnp/backend/extensions/vm/common.hpp index 4e19b5e12570..f2a46477d78c 100644 --- a/dpnp/backend/extensions/vm/common.hpp +++ b/dpnp/backend/extensions/vm/common.hpp @@ -365,7 +365,7 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, */ #define MACRO_POPULATE_DISPATCH_VECTORS(__name__) \ template \ - struct ContigFactory \ + struct __name__##_ContigFactory \ { \ fnT get() \ { \ @@ -380,7 +380,7 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, }; \ \ template \ - struct TypeMapFactory \ + struct __name__##_TypeMapFactory \ { \ std::enable_if_t::value, int> get() \ { \ @@ -391,10 +391,11 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, \ static void populate_dispatch_vectors(void) \ { \ - ext_ns::init_dispatch_vector( \ + ext_ns::init_dispatch_vector( \ output_typeid_vector); \ ext_ns::init_dispatch_vector(contig_dispatch_vector); \ + __name__##_ContigFactory>( \ + contig_dispatch_vector); \ }; /** @@ -404,7 +405,7 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, */ #define MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(__name__) \ template \ - struct ContigFactory \ + struct __name__##_ContigFactory \ { \ fnT get() \ { \ @@ -423,7 +424,7 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, }; \ \ template \ - struct TypeMapFactory \ + struct __name__##_TypeMapFactory \ { \ std::enable_if_t>::value, \ std::pair> \ @@ -438,10 +439,12 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, \ static void populate_dispatch_vectors(void) \ { \ - ext_ns::init_dispatch_vector, TypeMapFactory>( \ + ext_ns::init_dispatch_vector, \ + __name__##_TypeMapFactory>( \ output_typeid_vector); \ ext_ns::init_dispatch_vector(contig_dispatch_vector); \ + __name__##_ContigFactory>( \ + contig_dispatch_vector); \ }; /** @@ -451,7 +454,7 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, */ #define MACRO_POPULATE_DISPATCH_TABLES(__name__) \ template \ - struct ContigFactory \ + struct __name__##_ContigFactory \ { \ fnT get() \ { \ @@ -467,7 +470,7 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, }; \ \ template \ - struct TypeMapFactory \ + struct __name__##_TypeMapFactory \ { \ std::enable_if_t::value, int> get() \ { \ @@ -478,9 +481,10 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q, \ static void populate_dispatch_tables(void) \ { \ - ext_ns::init_dispatch_table( \ + ext_ns::init_dispatch_table( \ output_typeid_vector); \ ext_ns::init_dispatch_table(contig_dispatch_vector); \ + __name__##_ContigFactory>( \ + contig_dispatch_vector); \ }; } // namespace dpnp::extensions::vm::py_internal diff --git a/dpnp/backend/extensions/vm/erf_funcs.cpp b/dpnp/backend/extensions/vm/erf_funcs.cpp index 7ca642fe2fc5..b43b7a19945e 100644 --- a/dpnp/backend/extensions/vm/erf_funcs.cpp +++ b/dpnp/backend/extensions/vm/erf_funcs.cpp @@ -84,7 +84,7 @@ struct OutputType static int output_typeid_vector[td_ns::num_types]; template -struct TypeMapFactory +struct ErfTypeMapFactory { std::enable_if_t::value, int> get() { @@ -169,7 +169,7 @@ void init_erf_funcs(py::module_ m) using arrayT = dpnp::tensor::usm_ndarray; using impl::output_typeid_vector; - init_dispatch_vector(output_typeid_vector); + init_dispatch_vector(output_typeid_vector); auto erf_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src, const arrayT &dst) {