From b9ee15cd5360c3426ba0994bdce3f5b23b078ac6 Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" Date: Sun, 24 Mar 2024 01:09:12 +0000 Subject: [PATCH] Merge pull request #1400 from IntelPython/feature/move_experimental_decorators_to_stable Feature/move experimental decorators to stable b1865287dda1437a5047b109a8a5bcb43f2ffdb1 --- dev/.buildinfo | 2 +- dev/_sources/autoapi/index.rst.txt | 4 +- .../numba_dpex/core/boxing/index.rst.txt | 384 ++++ .../core/boxing/ranges/index.rst.txt | 57 + .../core/boxing/usm_ndarray/index.rst.txt | 58 + .../numba_dpex/core/caching/index.rst.txt | 231 +++ .../numba_dpex/core/compiler/index.rst.txt | 46 + .../numba_dpex/core/config/index.rst.txt | 95 + .../numba_dpex/core/datamodel/index.rst.txt | 7 + .../core/datamodel/models/index.rst.txt | 190 ++ .../numba_dpex/core/decorators/index.rst.txt | 230 +++ .../numba_dpex/core/descriptor/index.rst.txt | 178 ++ .../core/dpjit_dispatcher/index.rst.txt | 61 + .../numba_dpex/core/exceptions/index.rst.txt | 22 + .../autoapi/numba_dpex/core/index.rst.txt | 347 ++++ .../core/itanium_mangler/index.rst.txt | 201 ++ .../arg_pack_unpacker/index.rst.txt | 37 + .../kernel_interface/arrayobj/index.rst.txt | 60 + .../kernel_interface/dispatcher/index.rst.txt | 44 + .../core/kernel_interface/func/index.rst.txt | 181 ++ .../core/kernel_interface/index.rst.txt | 12 + .../kernel_base/index.rst.txt | 63 + .../kernel_interface/launcher/index.rst.txt | 51 + .../ranges_overloads/index.rst.txt | 24 + .../spirv_kernel/index.rst.txt | 71 + .../core/kernel_interface/utils/index.rst.txt | 66 + .../kernel_launcher}/index.rst.txt | 10 +- .../numba_dpex/core/parfors/index.rst.txt | 7 + .../core/parfors/kernel_builder/index.rst.txt | 79 + .../core/parfors/parfor_lowerer/index.rst.txt | 105 ++ .../parfors/reduction_helper/index.rst.txt | 65 + .../reduction_kernel_builder/index.rst.txt | 41 + .../core/passes/dufunc_inliner/index.rst.txt | 32 + .../numba_dpex/core/passes/index.rst.txt | 119 ++ .../parfor_legalize_cfd_pass/index.rst.txt | 122 ++ .../core/passes/passes/index.rst.txt | 164 ++ .../pipelines/dpjit_compiler/index.rst.txt | 82 + .../numba_dpex/core/pipelines/index.rst.txt | 7 + .../pipelines/kernel_compiler/index.rst.txt | 58 + .../core/runtime/context/index.rst.txt | 286 +++ .../numba_dpex/core/runtime/index.rst.txt | 7 + .../core/targets/dpjit_target/index.rst.txt | 155 ++ .../numba_dpex/core/targets/index.rst.txt | 7 + .../typeconv/array_conversion/index.rst.txt | 42 + .../numba_dpex/core/typeconv/index.rst.txt | 40 + .../core/types/array_type/index.rst.txt | 63 + .../core/types/dpctl_types/index.rst.txt | 124 ++ .../types/dpnp_ndarray_type/index.rst.txt | 39 + .../numba_dpex/core/types/index.rst.txt | 336 ++++ .../types/kernel_api/atomic_ref/index.rst.txt | 65 + .../core/types/kernel_api/index.rst.txt | 7 + .../kernel_api/index_space_ids/index.rst.txt | 119 ++ .../kernel_api/literal_intenum/index.rst.txt | 82 + .../kernel_api/local_accessor/index.rst.txt | 99 + .../types/kernel_api/ranges/index.rst.txt | 51 + .../numba_types_short_names/index.rst.txt | 87 + .../core/types/usm_ndarray_type/index.rst.txt | 74 + .../core/typing/dpnpdecl/index.rst.txt | 164 ++ .../numba_dpex/core/typing/index.rst.txt | 7 + .../core/typing/typeof/index.rst.txt | 112 ++ .../core/utils/caching_utils/index.rst.txt | 62 + .../numba_dpex/core/utils/index.rst.txt | 105 ++ .../index.rst.txt | 145 ++ .../core/utils/kernel_launcher/index.rst.txt | 254 +++ .../core/utils/kernel_templates/index.rst.txt | 61 + .../kernel_template_iface/index.rst.txt | 58 + .../range_kernel_template/index.rst.txt | 63 + .../reduction_template/index.rst.txt | 105 ++ .../core/utils/suai_helper/index.rst.txt | 63 + .../_atomic_fence_overloads/index.rst.txt | 44 - .../_atomic_ref_overloads/index.rst.txt | 193 -- .../_group_barrier_overloads/index.rst.txt | 55 - .../_index_space_id_overloads/index.rst.txt | 94 - .../_private_array_overloads/index.rst.txt | 52 - .../_registry/index.rst.txt | 32 - .../_spv_atomic_inst_helper/index.rst.txt | 74 - .../index.rst.txt | 9 - .../spv_fn_declarations/index.rst.txt | 81 - .../experimental/decorators/index.rst.txt | 69 - .../numba_dpex/experimental/index.rst.txt | 119 -- .../experimental/models/index.rst.txt | 97 - .../experimental/target/index.rst.txt | 146 -- .../experimental/testing/index.rst.txt | 39 - .../experimental/typeof/index.rst.txt | 103 -- .../experimental/types/index.rst.txt | 61 - .../kernel_api/atomic_ref/index.rst.txt | 2 +- .../numba_dpex/kernel_api/index.rst.txt | 2 +- dev/autoapi/index.html | 8 +- dev/autoapi/numba_dpex/core/boxing/index.html | 777 ++++++++ .../numba_dpex/core/boxing/ranges/index.html | 368 ++++ .../boxing/usm_ndarray}/index.html | 67 +- .../numba_dpex/core/caching/index.html | 585 ++++++ .../decorators => core/compiler}/index.html | 105 +- dev/autoapi/numba_dpex/core/config/index.html | 433 +++++ .../numba_dpex/core/datamodel/index.html | 291 +++ .../core/datamodel/models/index.html | 511 ++++++ .../numba_dpex/core/decorators/index.html | 557 ++++++ .../target => core/descriptor}/index.html | 197 +- .../dpjit_dispatcher}/index.html | 44 +- .../testing => core/exceptions}/index.html | 37 +- dev/autoapi/numba_dpex/core/index.html | 733 ++++++++ .../core/itanium_mangler/index.html | 550 ++++++ .../arg_pack_unpacker/index.html | 337 ++++ .../core/kernel_interface/arrayobj/index.html | 358 ++++ .../kernel_interface/dispatcher/index.html | 342 ++++ .../core/kernel_interface/func/index.html | 513 ++++++ .../kernel_interface}/index.html | 11 +- .../kernel_interface/kernel_base/index.html | 369 ++++ .../kernel_interface/launcher}/index.html | 52 +- .../ranges_overloads}/index.html | 25 +- .../kernel_interface/spirv_kernel/index.html | 379 ++++ .../core/kernel_interface/utils/index.html | 381 ++++ .../kernel_launcher}/index.html | 42 +- .../numba_dpex/core/parfors/index.html | 291 +++ .../core/parfors/kernel_builder/index.html | 397 ++++ .../core/parfors/parfor_lowerer/index.html | 417 +++++ .../core/parfors/reduction_helper/index.html | 369 ++++ .../reduction_kernel_builder/index.html | 346 ++++ .../core/passes/dufunc_inliner/index.html | 335 ++++ .../models => core/passes}/index.html | 141 +- .../parfor_legalize_cfd_pass/index.html | 429 +++++ .../passes/passes}/index.html | 212 ++- .../pipelines/dpjit_compiler}/index.html | 125 +- .../numba_dpex/core/pipelines/index.html | 291 +++ .../core/pipelines/kernel_compiler/index.html | 358 ++++ .../core/runtime/context/index.html | 640 +++++++ .../numba_dpex/core/runtime/index.html | 291 +++ .../core/targets/dpjit_target/index.html | 477 +++++ .../numba_dpex/core/targets/index.html | 291 +++ .../typeconv/array_conversion}/index.html | 44 +- .../numba_dpex/core/typeconv/index.html | 351 ++++ .../core/types/array_type/index.html | 368 ++++ .../core/types/dpctl_types/index.html | 449 +++++ .../core/types/dpnp_ndarray_type/index.html | 338 ++++ dev/autoapi/numba_dpex/core/types/index.html | 716 ++++++++ .../types/kernel_api/atomic_ref/index.html | 361 ++++ .../core/types/kernel_api/index.html | 291 +++ .../kernel_api/index_space_ids/index.html | 420 +++++ .../kernel_api/literal_intenum/index.html | 387 ++++ .../kernel_api/local_accessor/index.html | 405 +++++ .../core/types/kernel_api/ranges/index.html | 350 ++++ .../types/numba_types_short_names/index.html | 421 +++++ .../core/types/usm_ndarray_type/index.html | 378 ++++ .../core/typing/dpnpdecl/index.html | 492 +++++ dev/autoapi/numba_dpex/core/typing/index.html | 291 +++ .../typing/typeof}/index.html | 165 +- .../core/utils/caching_utils/index.html | 382 ++++ dev/autoapi/numba_dpex/core/utils/index.html | 439 +++++ .../kernel_flattened_args_builder/index.html | 455 +++++ .../core/utils/kernel_launcher/index.html | 613 +++++++ .../core/utils/kernel_templates/index.html | 368 ++++ .../kernel_template_iface/index.html | 364 ++++ .../range_kernel_template/index.html | 368 ++++ .../reduction_template/index.html | 414 +++++ .../core/utils/suai_helper/index.html | 374 ++++ .../_atomic_ref_overloads/index.html | 567 ------ .../numba_dpex/experimental/index.html | 438 ----- .../numba_dpex/experimental/typeof/index.html | 432 ----- .../kernel_api/atomic_fence/index.html | 4 +- .../kernel_api/atomic_ref/index.html | 6 +- .../numba_dpex/kernel_api/barrier/index.html | 4 +- .../kernel_api/flag_enum/index.html | 4 +- dev/autoapi/numba_dpex/kernel_api/index.html | 12 +- .../kernel_api/index_space_ids/index.html | 8 +- .../numba_dpex/kernel_api/launcher/index.html | 4 +- .../kernel_api/local_accessor/index.html | 4 +- .../kernel_api/memory_enums/index.html | 4 +- .../kernel_api/private_array/index.html | 4 +- .../numba_dpex/kernel_api/ranges/index.html | 4 +- dev/contribution_guide.html | 4 +- dev/examples.html | 4 +- dev/experimental/index.html | 10 +- dev/genindex.html | 1614 ++++++++++++++--- dev/getting_started.html | 4 +- dev/index.html | 4 +- dev/license.html | 4 +- dev/objects.inv | Bin 3646 -> 7583 bytes dev/overview.html | 4 +- dev/programming_model.html | 4 +- dev/py-modindex.html | 338 +++- dev/release-notes.html | 4 +- dev/search.html | 4 +- dev/searchindex.js | 2 +- dev/useful_links.html | 4 +- dev/user_guide/config.html | 4 +- dev/user_guide/debugging/altering.html | 4 +- dev/user_guide/debugging/backtrace.html | 12 +- dev/user_guide/debugging/breakpoints.html | 51 +- dev/user_guide/debugging/common_issues.html | 4 +- dev/user_guide/debugging/data.html | 4 +- .../debugging/debugging_environment.html | 43 +- dev/user_guide/debugging/features.html | 4 +- dev/user_guide/debugging/frame_info.html | 51 +- dev/user_guide/debugging/index.html | 43 +- dev/user_guide/debugging/limitations.html | 4 +- dev/user_guide/debugging/local_variables.html | 83 +- dev/user_guide/debugging/numba-0.55.html | 8 +- dev/user_guide/debugging/set_up_machine.html | 4 +- dev/user_guide/debugging/stepping.html | 51 +- dev/user_guide/debugging/symbols.html | 43 +- dev/user_guide/index.html | 4 +- .../kernel_programming/device-functions.html | 20 +- dev/user_guide/kernel_programming/index.html | 14 +- .../kernel_programming/reduction.html | 103 +- .../supported-python-features.html | 4 +- .../kernel_programming/synchronization.html | 4 +- 206 files changed, 31420 insertions(+), 3894 deletions(-) create mode 100644 dev/_sources/autoapi/numba_dpex/core/boxing/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/boxing/ranges/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/boxing/usm_ndarray/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/caching/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/compiler/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/config/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/datamodel/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/datamodel/models/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/decorators/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/descriptor/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/dpjit_dispatcher/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/exceptions/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/itanium_mangler/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/arg_pack_unpacker/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/arrayobj/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/dispatcher/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/func/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/kernel_base/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/launcher/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/ranges_overloads/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/spirv_kernel/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/kernel_interface/utils/index.rst.txt rename dev/_sources/autoapi/numba_dpex/{experimental/launcher => core/kernel_launcher}/index.rst.txt (88%) create mode 100644 dev/_sources/autoapi/numba_dpex/core/parfors/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/parfors/kernel_builder/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/parfors/parfor_lowerer/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/parfors/reduction_helper/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/parfors/reduction_kernel_builder/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/passes/dufunc_inliner/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/passes/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/passes/parfor_legalize_cfd_pass/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/passes/passes/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/pipelines/dpjit_compiler/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/pipelines/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/pipelines/kernel_compiler/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/runtime/context/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/runtime/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/targets/dpjit_target/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/targets/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/typeconv/array_conversion/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/typeconv/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/array_type/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/dpctl_types/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/dpnp_ndarray_type/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/kernel_api/atomic_ref/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index_space_ids/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/kernel_api/literal_intenum/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/kernel_api/local_accessor/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/kernel_api/ranges/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/numba_types_short_names/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/types/usm_ndarray_type/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/typing/dpnpdecl/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/typing/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/typing/typeof/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/caching_utils/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/kernel_flattened_args_builder/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/kernel_launcher/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/kernel_template_iface/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/range_kernel_template/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/reduction_template/index.rst.txt create mode 100644 dev/_sources/autoapi/numba_dpex/core/utils/suai_helper/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_fence_overloads/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_private_array_overloads/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_registry/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/spv_fn_declarations/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/decorators/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/models/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/target/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/testing/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/typeof/index.rst.txt delete mode 100644 dev/_sources/autoapi/numba_dpex/experimental/types/index.rst.txt create mode 100644 dev/autoapi/numba_dpex/core/boxing/index.html create mode 100644 dev/autoapi/numba_dpex/core/boxing/ranges/index.html rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads/_private_array_overloads => core/boxing/usm_ndarray}/index.html (77%) create mode 100644 dev/autoapi/numba_dpex/core/caching/index.html rename dev/autoapi/numba_dpex/{experimental/decorators => core/compiler}/index.html (72%) create mode 100644 dev/autoapi/numba_dpex/core/config/index.html create mode 100644 dev/autoapi/numba_dpex/core/datamodel/index.html create mode 100644 dev/autoapi/numba_dpex/core/datamodel/models/index.html create mode 100644 dev/autoapi/numba_dpex/core/decorators/index.html rename dev/autoapi/numba_dpex/{experimental/target => core/descriptor}/index.html (61%) rename dev/autoapi/numba_dpex/{experimental/types => core/dpjit_dispatcher}/index.html (83%) rename dev/autoapi/numba_dpex/{experimental/testing => core/exceptions}/index.html (87%) create mode 100644 dev/autoapi/numba_dpex/core/index.html create mode 100644 dev/autoapi/numba_dpex/core/itanium_mangler/index.html create mode 100644 dev/autoapi/numba_dpex/core/kernel_interface/arg_pack_unpacker/index.html create mode 100644 dev/autoapi/numba_dpex/core/kernel_interface/arrayobj/index.html create mode 100644 dev/autoapi/numba_dpex/core/kernel_interface/dispatcher/index.html create mode 100644 dev/autoapi/numba_dpex/core/kernel_interface/func/index.html rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads => core/kernel_interface}/index.html (96%) create mode 100644 dev/autoapi/numba_dpex/core/kernel_interface/kernel_base/index.html rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads => core/kernel_interface/launcher}/index.html (82%) rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads/_registry => core/kernel_interface/ranges_overloads}/index.html (89%) create mode 100644 dev/autoapi/numba_dpex/core/kernel_interface/spirv_kernel/index.html create mode 100644 dev/autoapi/numba_dpex/core/kernel_interface/utils/index.html rename dev/autoapi/numba_dpex/{experimental/launcher => core/kernel_launcher}/index.html (86%) create mode 100644 dev/autoapi/numba_dpex/core/parfors/index.html create mode 100644 dev/autoapi/numba_dpex/core/parfors/kernel_builder/index.html create mode 100644 dev/autoapi/numba_dpex/core/parfors/parfor_lowerer/index.html create mode 100644 dev/autoapi/numba_dpex/core/parfors/reduction_helper/index.html create mode 100644 dev/autoapi/numba_dpex/core/parfors/reduction_kernel_builder/index.html create mode 100644 dev/autoapi/numba_dpex/core/passes/dufunc_inliner/index.html rename dev/autoapi/numba_dpex/{experimental/models => core/passes}/index.html (70%) create mode 100644 dev/autoapi/numba_dpex/core/passes/parfor_legalize_cfd_pass/index.html rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads => core/passes/passes}/index.html (57%) rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper => core/pipelines/dpjit_compiler}/index.html (69%) create mode 100644 dev/autoapi/numba_dpex/core/pipelines/index.html create mode 100644 dev/autoapi/numba_dpex/core/pipelines/kernel_compiler/index.html create mode 100644 dev/autoapi/numba_dpex/core/runtime/context/index.html create mode 100644 dev/autoapi/numba_dpex/core/runtime/index.html create mode 100644 dev/autoapi/numba_dpex/core/targets/dpjit_target/index.html create mode 100644 dev/autoapi/numba_dpex/core/targets/index.html rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads/_atomic_fence_overloads => core/typeconv/array_conversion}/index.html (86%) create mode 100644 dev/autoapi/numba_dpex/core/typeconv/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/array_type/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/dpctl_types/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/dpnp_ndarray_type/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/kernel_api/atomic_ref/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/kernel_api/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/kernel_api/index_space_ids/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/kernel_api/literal_intenum/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/kernel_api/local_accessor/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/kernel_api/ranges/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/numba_types_short_names/index.html create mode 100644 dev/autoapi/numba_dpex/core/types/usm_ndarray_type/index.html create mode 100644 dev/autoapi/numba_dpex/core/typing/dpnpdecl/index.html create mode 100644 dev/autoapi/numba_dpex/core/typing/index.html rename dev/autoapi/numba_dpex/{experimental/_kernel_dpcpp_spirv_overloads/spv_fn_declarations => core/typing/typeof}/index.html (62%) create mode 100644 dev/autoapi/numba_dpex/core/utils/caching_utils/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/kernel_flattened_args_builder/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/kernel_launcher/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/kernel_templates/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/kernel_templates/kernel_template_iface/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/kernel_templates/range_kernel_template/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/kernel_templates/reduction_template/index.html create mode 100644 dev/autoapi/numba_dpex/core/utils/suai_helper/index.html delete mode 100644 dev/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads/index.html delete mode 100644 dev/autoapi/numba_dpex/experimental/index.html delete mode 100644 dev/autoapi/numba_dpex/experimental/typeof/index.html diff --git a/dev/.buildinfo b/dev/.buildinfo index 5ec862cc00..01471813a6 100644 --- a/dev/.buildinfo +++ b/dev/.buildinfo @@ -1,4 +1,4 @@ # Sphinx build info version 1 # This file hashes the configuration used when building these files. When it is not found, a full rebuild will be done. -config: 35f385f41b2d9b7b4ed069614fa92a39 +config: de748025fac40afc5045ed95f285b3df tags: 645f666f9bcd5a90fca523b33c5a78b7 diff --git a/dev/_sources/autoapi/index.rst.txt b/dev/_sources/autoapi/index.rst.txt index 0745c4161b..056d157624 100644 --- a/dev/_sources/autoapi/index.rst.txt +++ b/dev/_sources/autoapi/index.rst.txt @@ -7,8 +7,8 @@ This page contains auto-generated API reference documentation [#f1]_. :maxdepth: 1 numba_dpex/kernel_api/index - numba_dpex/experimental/decorators/index - numba_dpex/experimental/launcher/index + numba_dpex/core/decorators/index + numba_dpex/core/kernel_launcher/index .. [#f1] Created with `sphinx-autoapi `_ \ No newline at end of file diff --git a/dev/_sources/autoapi/numba_dpex/core/boxing/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/boxing/index.rst.txt new file mode 100644 index 0000000000..18582d648e --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/boxing/index.rst.txt @@ -0,0 +1,384 @@ + +numba_dpex.core.boxing +====================== + +.. py:module:: numba_dpex.core.boxing + +.. autoapi-nested-parse:: + + Contains the ``box`` and ``unbox`` functions for numba_dpex types that are + passable as arguments to a kernel or dpjit decorated function. + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`NdRangeType ` + - Numba-dpex type corresponding to + * - :py:obj:`RangeType ` + - Numba-dpex type corresponding to + * - :py:obj:`NdRange ` + - A class to encapsulate all kernel launch parameters. + * - :py:obj:`Range ` + - A data structure to encapsulate a single kernel launch parameter. + * - :py:obj:`USMNdArray ` + - A type class to represent dpctl.tensor.usm_ndarray. + * - :py:obj:`NdRange ` + - A class to encapsulate all kernel launch parameters. + * - :py:obj:`Range ` + - A data structure to encapsulate a single kernel launch parameter. + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`unbox_range `\ (typ, obj, c) + - Converts a Python Range object to numba-dpex's native struct representation + * - :py:obj:`unbox_ndrange `\ (typ, obj, c) + - Converts a Python Range object to numba-dpex's native struct representation + * - :py:obj:`box_range `\ (typ, val, c) + - Convert a native range structure to a Range object. + * - :py:obj:`box_ndrange `\ (typ, val, c) + - Convert a native range structure to a Range object. + * - :py:obj:`unbox_dpnp_nd_array `\ (typ, obj, c) + - Converts a dpctl.tensor.usm_ndarray/dpnp.ndarray object to a Numba-dpex + * - :py:obj:`box_array `\ (typ, val, c) + - Boxes a NativeValue representation of USMNdArray/DpnpNdArray type into a + + + +Classes +------- + +.. py:class:: NdRangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.NdRange` + + + + +.. py:class:: RangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.Range` + + + + +.. py:class:: NdRange(global_size, local_size) + + A class to encapsulate all kernel launch parameters. + + The NdRange defines the index space for a work group as well as + the global index space. It is passed to parallel_for to execute + a kernel on a set of work items. + + This class basically contains two Range object, one for the global_range + and the other for the local_range. The global_range parameter contains + the global index space and the local_range parameter contains the index + space of a work group. This class mimics the behavior of `sycl::nd_range` + class. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get_global_range `\ () + - Returns a Range defining the index space. + * - :py:obj:`get_local_range `\ () + - Returns a Range defining the index space of a work group. + + + .. rubric:: Members + + .. py:method:: get_global_range() + + Returns a Range defining the index space. + + :returns: A `Range` object defining the index space. + :rtype: Range + + + .. py:method:: get_local_range() + + Returns a Range defining the index space of a work group. + + :returns: A `Range` object to specify index space of a work group. + :rtype: Range + + + + +.. py:class:: Range + + Bases: :py:obj:`tuple` + + A data structure to encapsulate a single kernel launch parameter. + + The range is an abstraction that describes the number of elements + in each dimension of buffers and index spaces. It can contain + 1, 2, or 3 numbers, depending on the dimensionality of the + object it describes. + + This is just a wrapper class on top of a 3-tuple. The kernel launch + parameter is consisted of three int's. This class basically mimics + the behavior of `sycl::range`. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get `\ (index) + - Returns the range of a single dimension. + * - :py:obj:`size `\ () + - Returns the size of a range. + + + .. rubric:: Members + + .. py:method:: get(index) + + Returns the range of a single dimension. + + :param index: The index of the dimension, i.e. [0,2] + :type index: int + + :returns: The range of the dimension indexed by `index`. + :rtype: int + + + .. py:method:: size() + + Returns the size of a range. + + Returns the size of a range by multiplying + the range of the individual dimensions. + + :returns: The size of a range. + :rtype: int + + + + +.. py:class:: USMNdArray(ndim, layout='C', dtype=None, usm_type='device', device=None, queue=None, readonly=False, name=None, aligned=True, addrspace=address_space.GLOBAL) + + Bases: :py:obj:`numba.core.types.npytypes.Array` + + A type class to represent dpctl.tensor.usm_ndarray. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy `\ (dtype, ndim, layout, readonly, addrspace, device, usm_type) + - \- + * - :py:obj:`unify `\ (typingctx, other) + - Unify this with the *other* USMNdArray. + * - :py:obj:`can_convert_to `\ (typingctx, other) + - Convert this USMNdArray to the *other*. + + + .. rubric:: Members + + .. py:method:: copy(dtype=None, ndim=None, layout=None, readonly=None, addrspace=None, device=None, usm_type=None) + + + .. py:method:: unify(typingctx, other) + + Unify this with the *other* USMNdArray. + + + .. py:method:: can_convert_to(typingctx, other) + + Convert this USMNdArray to the *other*. + + + + +.. py:class:: NdRange(global_size, local_size) + + A class to encapsulate all kernel launch parameters. + + The NdRange defines the index space for a work group as well as + the global index space. It is passed to parallel_for to execute + a kernel on a set of work items. + + This class basically contains two Range object, one for the global_range + and the other for the local_range. The global_range parameter contains + the global index space and the local_range parameter contains the index + space of a work group. This class mimics the behavior of `sycl::nd_range` + class. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get_global_range `\ () + - Returns a Range defining the index space. + * - :py:obj:`get_local_range `\ () + - Returns a Range defining the index space of a work group. + + + .. rubric:: Members + + .. py:method:: get_global_range() + + Returns a Range defining the index space. + + :returns: A `Range` object defining the index space. + :rtype: Range + + + .. py:method:: get_local_range() + + Returns a Range defining the index space of a work group. + + :returns: A `Range` object to specify index space of a work group. + :rtype: Range + + + + +.. py:class:: Range + + Bases: :py:obj:`tuple` + + A data structure to encapsulate a single kernel launch parameter. + + The range is an abstraction that describes the number of elements + in each dimension of buffers and index spaces. It can contain + 1, 2, or 3 numbers, depending on the dimensionality of the + object it describes. + + This is just a wrapper class on top of a 3-tuple. The kernel launch + parameter is consisted of three int's. This class basically mimics + the behavior of `sycl::range`. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get `\ (index) + - Returns the range of a single dimension. + * - :py:obj:`size `\ () + - Returns the size of a range. + + + .. rubric:: Members + + .. py:method:: get(index) + + Returns the range of a single dimension. + + :param index: The index of the dimension, i.e. [0,2] + :type index: int + + :returns: The range of the dimension indexed by `index`. + :rtype: int + + + .. py:method:: size() + + Returns the size of a range. + + Returns the size of a range by multiplying + the range of the individual dimensions. + + :returns: The size of a range. + :rtype: int + + + + +Functions +--------- +.. py:function:: unbox_range(typ, obj, c) + + Converts a Python Range object to numba-dpex's native struct representation + for RangeType. + + +.. py:function:: unbox_ndrange(typ, obj, c) + + Converts a Python Range object to numba-dpex's native struct representation + for NdRangeType. + + +.. py:function:: box_range(typ, val, c) + + Convert a native range structure to a Range object. + + +.. py:function:: box_ndrange(typ, val, c) + + Convert a native range structure to a Range object. + + +.. py:function:: unbox_dpnp_nd_array(typ, obj, c) + + Converts a dpctl.tensor.usm_ndarray/dpnp.ndarray object to a Numba-dpex + internal array structure. + + :param typ: The Numba type of the PyObject + :param obj: The actual PyObject to be unboxed + :param c: The unboxing context + + Returns: A NativeValue object representing an unboxed + dpctl.tensor.usm_ndarray/dpnp.ndarray + + +.. py:function:: box_array(typ, val, c) + + Boxes a NativeValue representation of USMNdArray/DpnpNdArray type into a + dpctl.tensor.usm_ndarray/dpnp.ndarray PyObject + + :param typ: The representation of the USMNdArray/DpnpNdArray type. + :param val: A native representation of a Numba USMNdArray/DpnpNdArray type + object. + :param c: The boxing context. + + Returns: A Pyobject for a dpctl.tensor.usm_ndarray/dpnp.ndarray boxed from + the Numba-dpex native value. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/boxing/ranges/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/boxing/ranges/index.rst.txt new file mode 100644 index 0000000000..0d85a6e862 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/boxing/ranges/index.rst.txt @@ -0,0 +1,57 @@ + +:orphan: + +numba_dpex.core.boxing.ranges +============================= + +.. py:module:: numba_dpex.core.boxing.ranges + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`unbox_range `\ (typ, obj, c) + - Converts a Python Range object to numba-dpex's native struct representation + * - :py:obj:`unbox_ndrange `\ (typ, obj, c) + - Converts a Python Range object to numba-dpex's native struct representation + * - :py:obj:`box_range `\ (typ, val, c) + - Convert a native range structure to a Range object. + * - :py:obj:`box_ndrange `\ (typ, val, c) + - Convert a native range structure to a Range object. + + + + +Functions +--------- +.. py:function:: unbox_range(typ, obj, c) + + Converts a Python Range object to numba-dpex's native struct representation + for RangeType. + + +.. py:function:: unbox_ndrange(typ, obj, c) + + Converts a Python Range object to numba-dpex's native struct representation + for NdRangeType. + + +.. py:function:: box_range(typ, val, c) + + Convert a native range structure to a Range object. + + +.. py:function:: box_ndrange(typ, val, c) + + Convert a native range structure to a Range object. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/boxing/usm_ndarray/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/boxing/usm_ndarray/index.rst.txt new file mode 100644 index 0000000000..707ccd933b --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/boxing/usm_ndarray/index.rst.txt @@ -0,0 +1,58 @@ + +:orphan: + +numba_dpex.core.boxing.usm_ndarray +================================== + +.. py:module:: numba_dpex.core.boxing.usm_ndarray + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`unbox_dpnp_nd_array `\ (typ, obj, c) + - Converts a dpctl.tensor.usm_ndarray/dpnp.ndarray object to a Numba-dpex + * - :py:obj:`box_array `\ (typ, val, c) + - Boxes a NativeValue representation of USMNdArray/DpnpNdArray type into a + + + + +Functions +--------- +.. py:function:: unbox_dpnp_nd_array(typ, obj, c) + + Converts a dpctl.tensor.usm_ndarray/dpnp.ndarray object to a Numba-dpex + internal array structure. + + :param typ: The Numba type of the PyObject + :param obj: The actual PyObject to be unboxed + :param c: The unboxing context + + Returns: A NativeValue object representing an unboxed + dpctl.tensor.usm_ndarray/dpnp.ndarray + + +.. py:function:: box_array(typ, val, c) + + Boxes a NativeValue representation of USMNdArray/DpnpNdArray type into a + dpctl.tensor.usm_ndarray/dpnp.ndarray PyObject + + :param typ: The representation of the USMNdArray/DpnpNdArray type. + :param val: A native representation of a Numba USMNdArray/DpnpNdArray type + object. + :param c: The boxing context. + + Returns: A Pyobject for a dpctl.tensor.usm_ndarray/dpnp.ndarray boxed from + the Numba-dpex native value. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/caching/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/caching/index.rst.txt new file mode 100644 index 0000000000..c8c72ea927 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/caching/index.rst.txt @@ -0,0 +1,231 @@ + +:orphan: + +numba_dpex.core.caching +======================= + +.. py:module:: numba_dpex.core.caching + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`AbstractCache ` + - Abstract cache class to specify basic caching operations. + * - :py:obj:`NullCache ` + - A dummy cache used if user decides to disable caching. + * - :py:obj:`Node ` + - A 'Node' class for LRUCache. + * - :py:obj:`LRUCache ` + - LRUCache implementation for caching kernels, + + + + +Classes +------- + +.. py:class:: AbstractCache + + Abstract cache class to specify basic caching operations. + + This class will be used to create an non-functional dummy cache + (i.e. NullCache) and other functional cache. The dummy cache + will be used as a placeholder when caching is disabled. + + :param metaclass: Metaclass for the abstract class. + Defaults to ABCMeta. + :type metaclass: type, optional + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get `\ () + - :summarylabel:`abc` An abstract method to retrieve item from the cache. + * - :py:obj:`put `\ (key, value) + - :summarylabel:`abc` An abstract method to save item into the cache. + + + .. rubric:: Members + + .. py:method:: get() + :abstractmethod: + + An abstract method to retrieve item from the cache. + + + .. py:method:: put(key, value) + :abstractmethod: + + An abstract method to save item into the cache. + + :param key: The key for the data + (i.e. compiled kernel/function etc.). + :type key: object + :param value: The data (i.e. compiled kernel/function) + to be saved. + :type value: object + + + + +.. py:class:: NullCache + + Bases: :py:obj:`AbstractCache` + + A dummy cache used if user decides to disable caching. + + If the caching is disabled this class will be used to + perform all caching operations, all of which will be basically + NOP. This idea is copied from numba. + + :param AbstractCache: The abstract cache from which all + :type AbstractCache: class + :param other caching classes will be derived.: + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get `\ (key) + - Function to get an item (i.e. compiled kernel/function) + * - :py:obj:`put `\ (key, value) + - Function to save a compiled kernel/function + + + .. rubric:: Members + + .. py:method:: get(key) + + Function to get an item (i.e. compiled kernel/function) + from the cache + + :param key: The key to retrieve the + data (i.e. compiled kernel/function) + :type key: object + + :returns: Returns None. + :rtype: None + + + .. py:method:: put(key, value) + + Function to save a compiled kernel/function + into the cache. + + :param key: The key to the data (i.e. compiled kernel/function). + :type key: object + :param value: The data to be cached (i.e. + :type value: object + :param compiled kernel/function).: + + + + +.. py:class:: Node(key, value) + + A 'Node' class for LRUCache. + + + + +.. py:class:: LRUCache(name='cache', capacity=10, pyfunc=None) + + Bases: :py:obj:`AbstractCache` + + LRUCache implementation for caching kernels, + functions and modules. + + The cache is basically a doubly-linked-list backed + with a dictionary as a lookup table. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`size `\ () + - Get the current size of the cache. + * - :py:obj:`memsize `\ () + - Get the total memory size of the cache. + * - :py:obj:`clean `\ () + - Clean the cache + * - :py:obj:`get `\ (key) + - Get the value associated with the key. + * - :py:obj:`put `\ (key, value) + - Store the key-value pair into the cache. + + + .. rubric:: Members + + .. py:method:: size() + + Get the current size of the cache. + + :returns: The current number of items in the cache. + :rtype: int + + + .. py:method:: memsize() + + Get the total memory size of the cache. + + This function might be useful in the future when + size based (not count based) cache limit will be + implemented. + + :returns: Get the total memory size of the cache in bytes. + :rtype: int + + + .. py:method:: clean() + + Clean the cache + + + .. py:method:: get(key) + + Get the value associated with the key. + + :param key: A key for the lookup table. + :type key: object + + :returns: The value associated with the key. + :rtype: object + + + .. py:method:: put(key, value) + + Store the key-value pair into the cache. + + :param key: The key for the data. + :type key: object + :param value: The data to be saved. + :type value: object + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/compiler/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/compiler/index.rst.txt new file mode 100644 index 0000000000..ca92c59561 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/compiler/index.rst.txt @@ -0,0 +1,46 @@ + +:orphan: + +numba_dpex.core.compiler +======================== + +.. py:module:: numba_dpex.core.compiler + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`compile_with_dpex `\ (pyfunc, pyfunc_name, args, return_type, target_context, typing_context, debug, is_kernel, extra_compile_flags) + - Compiles a function using class:`numba_dpex.core.pipelines.KernelCompiler` + + + + +Functions +--------- +.. py:function:: compile_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. + + :param args: The list of arguments passed to the kernel. + :param debug: Optional flag to turn on debug mode compilation. + :type debug: bool + :param extra_compile_flags: Extra flags passed to the compiler. + + :returns: Compiled result. + :rtype: cres + + :raises KernelHasReturnValueError: If the compiled function returns a + :raises non-void value.: + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/config/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/config/index.rst.txt new file mode 100644 index 0000000000..ca9449a700 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/config/index.rst.txt @@ -0,0 +1,95 @@ + +:orphan: + +numba_dpex.core.config +====================== + +.. py:module:: numba_dpex.core.config + + +Overview +-------- + + + + +Attributes +---------- +.. py:data:: SAVE_IR_FILES + + + +.. py:data:: OFFLOAD_DIAGNOSTICS + + + +.. py:data:: DEBUG + + + +.. py:data:: DEBUGINFO_DEFAULT + + + +.. py:data:: DUMP_KERNEL_LLVM + + + +.. py:data:: DUMP_KERNEL_LAUNCHER + + + +.. py:data:: DEBUG_KERNEL_LAUNCHER + + + +.. py:data:: BUILD_KERNEL_OPTIONS + + + +.. py:data:: ENABLE_CACHE + + + +.. py:data:: CACHE_SIZE + + + +.. py:data:: DEBUG_CACHE + + + +.. py:data:: STATIC_LOCAL_MEM_PASS + + + +.. py:data:: TESTING_SKIP_NO_DPNP + + + +.. py:data:: TESTING_SKIP_NO_DEBUGGING + + + +.. py:data:: TESTING_LOG_DEBUGGING + + + +.. py:data:: STATIC_LOCAL_MEM_PASS + + + +.. py:data:: DPEX_OPT + + + +.. py:data:: INLINE_THRESHOLD + + + +.. py:data:: USE_MLIR + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/datamodel/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/datamodel/index.rst.txt new file mode 100644 index 0000000000..a4b03c4e7e --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/datamodel/index.rst.txt @@ -0,0 +1,7 @@ + +numba_dpex.core.datamodel +========================= + +.. py:module:: numba_dpex.core.datamodel + + diff --git a/dev/_sources/autoapi/numba_dpex/core/datamodel/models/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/datamodel/models/index.rst.txt new file mode 100644 index 0000000000..ef3e683e03 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/datamodel/models/index.rst.txt @@ -0,0 +1,190 @@ + +:orphan: + +numba_dpex.core.datamodel.models +================================ + +.. py:module:: numba_dpex.core.datamodel.models + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`GenericPointerModel ` + - A primitive type can be represented natively in the target in all + * - :py:obj:`IntEnumLiteralModel ` + - Representation of an object of LiteralIntEnum type using Numba's + * - :py:obj:`USMArrayDeviceModel ` + - A data model to represent a usm array type in the LLVM IR generated for a + * - :py:obj:`USMArrayHostModel ` + - Data model for the USMNdArray type when used in a host-only function. + * - :py:obj:`SyclQueueModel ` + - Represents the native data model for a dpctl.SyclQueue PyObject. + * - :py:obj:`SyclEventModel ` + - Represents the native data model for a dpctl.SyclEvent PyObject. + * - :py:obj:`RangeModel ` + - The numba_dpex data model for a numba_dpex.kernel_api.Range PyObject. + * - :py:obj:`NdRangeModel ` + - The numba_dpex data model for a numba_dpex.kernel_api.NdRange PyObject. + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get_flattened_member_count `\ (ty) + - Returns the number of fields in an instance of a given StructModel. + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dpex_data_model_manager ` + - \- + + +Classes +------- + +.. py:class:: GenericPointerModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.PrimitiveModel` + + A primitive type can be represented natively in the target in all + usage contexts. + + + + +.. py:class:: IntEnumLiteralModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.PrimitiveModel` + + Representation of an object of LiteralIntEnum type using Numba's + PrimitiveModel that can be represented natively in the target in all + usage contexts. + + + + +.. py:class:: USMArrayDeviceModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.StructModel` + + A data model to represent a usm array type in the LLVM IR generated for a + device-only kernel function. + + The USMArrayDeviceModel adds an extra address space attribute to the data + member. The extra attribute is needed when passing usm_ndarray array + arguments to kernels that are compiled for certain OpenCL GPU devices. Note + that the address space attribute is applied only to the data member and not + other members of USMArrayDeviceModel that are pointers. It is done this way + as other pointer members such as meminfo are not used inside a kernel and + these members maybe removed from the USMArrayDeviceModel in + future (refer #929). + + We use separate data models for host (USMArrayHostModel) and device + (USMArrayDeviceModel) as the address space attribute is only required for + kernel functions and not needed for functions that are compiled for a host + memory space. + + + + +.. py:class:: USMArrayHostModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.StructModel` + + Data model for the USMNdArray type when used in a host-only function. + + USMArrayHostModel is used by the numba_dpex.types.USMNdArray and + numba_dpex.types.DpnpNdArray type and abstracts the usmarystruct_t C type + defined in numba_dpex.core.runtime._usmarraystruct.h. + + The USMArrayDeviceModel differs from numba's ArrayModel by including an + extra member sycl_queue that maps to _usmarraystruct.sycl_queue pointer. The + _usmarraystruct.sycl_queue pointer stores the C++ sycl::queue pointer that + was used to allocate the data for the dpctl.tensor.usm_ndarray or + dpnp.ndarray represented by an instance of _usmarraystruct. + + + + +.. py:class:: SyclQueueModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.StructModel` + + Represents the native data model for a dpctl.SyclQueue PyObject. + + Numba-dpex uses a C struct as defined in + numba_dpex/core/runtime._queuestruct.h to store the required attributes for + a ``dpctl.SyclQueue`` Python object. + + - ``queue_ref``: An opaque C pointer to an actual SYCL queue C++ object. + - ``parent``: A PyObject* that stores a reference back to the original + ``dpctl.SyclQueue`` PyObject if the native struct is + created by unboxing the PyObject. + + + + +.. py:class:: SyclEventModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.StructModel` + + Represents the native data model for a dpctl.SyclEvent PyObject. + + Numba-dpex uses a C struct as defined in + numba_dpex/core/runtime._eventstruct.h to store the required attributes for + a ``dpctl.SyclEvent`` Python object. + + - ``event_ref``: An opaque C pointer to an actual SYCL event C++ object. + - ``parent``: A PyObject* that stores a reference back to the original + ``dpctl.SyclEvent`` PyObject if the native struct is + created by unboxing the PyObject. + + + + +.. py:class:: RangeModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.StructModel` + + The numba_dpex data model for a numba_dpex.kernel_api.Range PyObject. + + + + +.. py:class:: NdRangeModel(dmm, fe_type) + + Bases: :py:obj:`numba.core.datamodel.models.StructModel` + + The numba_dpex data model for a numba_dpex.kernel_api.NdRange PyObject. + + + + +Functions +--------- +.. py:function:: get_flattened_member_count(ty) + + Returns the number of fields in an instance of a given StructModel. + + + +Attributes +---------- +.. py:data:: dpex_data_model_manager + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/decorators/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/decorators/index.rst.txt new file mode 100644 index 0000000000..e5c9005097 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/decorators/index.rst.txt @@ -0,0 +1,230 @@ + +:orphan: + +numba_dpex.core.decorators +========================== + +.. py:module:: numba_dpex.core.decorators + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`kernel `\ (function_or_signature, \*\*options) + - A decorator to compile a function written using :py:mod:`numba_dpex.kernel_api`. + * - :py:obj:`device_func `\ (function_or_signature, \*\*options) + - Compiles a device-callable function that can be only invoked from a kernel. + * - :py:obj:`dpjit `\ (\*args, \*\*kws) + - \- + + + + +Functions +--------- +.. py:function:: kernel(function_or_signature=None, **options) + + A decorator to compile a function written using :py:mod:`numba_dpex.kernel_api`. + + The ``kernel`` decorator triggers the compilation of a function written + using the data-parallel kernel programming API exposed by + :py:mod:`numba_dpex.kernel_api`. Such a function is conceptually + equivalent to a kernel function written in the C++ SYCL eDSL. The + decorator will compile the function based on the types of the arguments + to a SPIR-V binary that can be executed either on OpenCL CPU, GPU + devices or Intel Level Zero GPU devices. + + Any function to be compilable using the kernel decorator should + adhere to the following semantic rules: + + - The first argument to the function should be either an instance of the + :class:`numba_dpex.kernel_api.Item` class or an instance of the + :class:`numba_dpex.kernel_api.NdItem`. + + - The function should not return any value. + + - The function should have at least one array type argument that can + either be an instance of ``dpnp.ndarray`` or an instance of + ``dpctl.tensor.usm_ndarray``. + + + :param signature_or_function: An optional signature or list of + signatures for which a function is to be compiled. Passing in a + signature "specializes" the decorated function and no other versions + of the function will be compiled. A function can also be + directly passed instead of a signature and the signature will get + inferred from the function. The actual compilation happens on every + invocation of the :func:`numba_dpex.experimental.call_kernel` + function where the decorated function is passed in as an argument + along with the argument values for the decorated function. + :type signature_or_function: optional + :param options: + - **debug** (bool): Whether the compilation should happen in debug + mode. *(Default = False)* + - **inline_threshold** (int): Specifies the level of inlining that + the compiler should attempt. *(Default = 2)* + :type options: optional + + :returns: An instance of + :class:`numba_dpex.kernel_api_impl.spirv.dispatcher.KernelDispatcher`. + The ``KernelDispatcher`` object compiles the decorated function when + passed in to :func:`numba_dpex.experimental.call_kernel`. + + Examples: + + 1. Decorate a function and pass it to ``call_kernel`` for compilation and + execution. + + .. code-block:: python + + import dpnp + import numba_dpex as dpex + from numba_dpex import kernel_api as kapi + + + # Data parallel kernel implementing vector sum + @dpex.kernel + def vecadd(item: kapi.Item, a, b, c): + i = item.get_id(0) + c[i] = a[i] + b[i] + + + N = 1024 + a = dpnp.ones(N) + b = dpnp.ones_like(a) + c = dpnp.zeros_like(a) + dpex.call_kernel(vecadd, kapi.Range(N), a, b, c) + + 2. Specializes a kernel and then compiles it directly before executing it + via ``call_kernel``. The kernel is specialized to expect a 1-D + ``dpnp.ndarray`` with either ``float32`` type elements or ``int64`` type + elements. + + .. code-block:: python + + import dpnp + import numba_dpex as dpex + from numba_dpex import kernel_api as kapi + from numba_dpex import DpnpNdArray, float32, int64 + from numba_dpex.core.types.kernel_api.index_space_ids import ItemType + + i64arrty = DpnpNdArray(ndim=1, dtype=int64, layout="C") + f32arrty = DpnpNdArray(ndim=1, dtype=float32, layout="C") + item_ty = ItemType(ndim=1) + + specialized_kernel = dpex.kernel( + [ + (item_ty, i64arrty, i64arrty, i64arrty), + (item_ty, f32arrty, f32arrty, f32arrty), + ] + ) + + + def vecadd(item: kapi.Item, a, b, c): + i = item.get_id(0) + c[i] = a[i] + b[i] + + + # Compile all specializations for vecadd + precompiled_kernels = specialized_kernel(vecadd) + N = 1024 + a = dpnp.ones(N, dtype=dpnp.int64) + b = dpnp.ones_like(a) + c = dpnp.zeros_like(a) + # Call a specific pre-compiled version of vecadd + dpex.call_kernel(precompiled_kernels, kapi.Range(N), a, b, c) + + + +.. py:function:: device_func(function_or_signature=None, **options) + + Compiles a device-callable function that can be only invoked from a kernel. + + The decorator is used to express auxiliary device-only functions that can + be called from a kernel or another device function, but are not callable + from the host. This decorator :func:`numba_dpex.experimental.device_func` + has no direct analogue in SYCL and primarily is provided to help programmers + make their kapi applications modular. + + A ``device_func`` decorated function does not require the first argument to + be a :class:`numba_dpex.kernel_api.Item` object or a + :class:`numba_dpex.kernel_api.NdItem` object, and unlike a ``kernel`` + decorated function is allowed to return any value. + All :py:mod:`numba_dpex.kernel_api` functionality can be used in a + ``device_func`` decorated function. + + The decorator is also used to compile overloads in the ``DpexKernelTarget``. + + A ``device_func`` decorated function is not compiled down to device binary + and instead is compiled down to LLVM IR. Final compilation to binary happens + when the function is invoked from a ``kernel`` decorated function. The + compilation happens this was to allow a ``device_func`` decorated function + to be internally linked into the kernel module at the LLVM level, leading to + more optimization opportunities. + + :param signature_or_function: An optional signature or list of + signatures for which a function is to be compiled. Passing in a + signature "specializes" the decorated function and no other versions + of the function will be compiled. A function can also be + directly passed instead of a signature and the signature will get + inferred from the function. The actual compilation happens on every + invocation of the decorated function from another ``device_func`` or + ``kernel`` decorated function. + :type signature_or_function: optional + :param options: + - **debug** (bool): Whether the compilation should happen in debug + mode. *(Default = False)* + - **inline_threshold** (int): Specifies the level of inlining that + the compiler should attempt. *(Default = 2)* + :type options: optional + + :returns: An instance of + :class:`numba_dpex.kernel_api_impl.spirv.dispatcher.KernelDispatcher`. + The ``KernelDispatcher`` object compiles the decorated function when + it is called from another function. + + Example: + + .. code-block:: python + + import dpnp + + from numba_dpex import experimental as dpex_exp + from numba_dpex import kernel_api as kapi + + + @dpex_exp.device_func + def increment_value(nd_item: NdItem, a): + i = nd_item.get_global_id(0) + + a[i] += 1 + group_barrier(nd_item.get_group(), MemoryScope.DEVICE) + + if i == 0: + for idx in range(1, a.size): + a[0] += a[idx] + + + @dpex_exp.kernel + def another_kernel(nd_item: NdItem, a): + increment_value(nd_item, a) + + + N = 16 + b = dpnp.ones(N, dtype=dpnp.int32) + + dpex_exp.call_kernel(another_kernel, dpex.NdRange((N,), (N,)), b) + + +.. py:function:: dpjit(*args, **kws) + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/descriptor/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/descriptor/index.rst.txt new file mode 100644 index 0000000000..65ee156ae3 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/descriptor/index.rst.txt @@ -0,0 +1,178 @@ + +:orphan: + +numba_dpex.core.descriptor +========================== + +.. py:module:: numba_dpex.core.descriptor + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpexTargetOptions ` + - Target options maps user options from decorators to the + * - :py:obj:`DpexKernelTarget ` + - Implements a target descriptor for numba_dpex.kernel decorated functions. + * - :py:obj:`DpexTarget ` + - Implements a target descriptor for numba_dpex.dpjit decorated functions. + + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dpex_kernel_target ` + - \- + * - :py:obj:`dpex_target ` + - \- + + +Classes +------- + +.. py:class:: DpexTargetOptions + + Bases: :py:obj:`numba.core.cpu.CPUTargetOptions` + + Target options maps user options from decorators to the + ``numba.core.compiler.Flags`` used by lowering and target context. + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`experimental ` + - \- + * - :py:obj:`release_gil ` + - \- + * - :py:obj:`no_compile ` + - \- + * - :py:obj:`use_mlir ` + - \- + * - :py:obj:`inline_threshold ` + - \- + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`finalize `\ (flags, options) + - Subclasses can override this method to make target specific + + + .. rubric:: Members + + .. py:attribute:: experimental + + + + .. py:attribute:: release_gil + + + + .. py:attribute:: no_compile + + + + .. py:attribute:: use_mlir + + + + .. py:attribute:: inline_threshold + + + + .. py:method:: finalize(flags, options) + + Subclasses can override this method to make target specific + customizations of default flags. + + :param flags: + :type flags: Flags + :param options: + :type options: dict + + + + +.. py:class:: DpexKernelTarget(target_name) + + Bases: :py:obj:`numba.core.descriptors.TargetDescriptor` + + Implements a target descriptor for numba_dpex.kernel decorated functions. + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`options ` + - \- + + + + .. rubric:: Members + + .. py:attribute:: options + + + + + +.. py:class:: DpexTarget(target_name) + + Bases: :py:obj:`numba.core.descriptors.TargetDescriptor` + + Implements a target descriptor for numba_dpex.dpjit decorated functions. + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`options ` + - \- + + + + .. rubric:: Members + + .. py:attribute:: options + + + + + + +Attributes +---------- +.. py:data:: dpex_kernel_target + + + +.. py:data:: dpex_target + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/dpjit_dispatcher/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/dpjit_dispatcher/index.rst.txt new file mode 100644 index 0000000000..e056518ae4 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/dpjit_dispatcher/index.rst.txt @@ -0,0 +1,61 @@ + +:orphan: + +numba_dpex.core.dpjit_dispatcher +================================ + +.. py:module:: numba_dpex.core.dpjit_dispatcher + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpjitDispatcher ` + - A dpex.djit-specific dispatcher. + + + + +Classes +------- + +.. py:class:: DpjitDispatcher(py_func, locals={}, targetoptions={}, pipeline_class=dpjit_compiler.DpjitCompiler) + + Bases: :py:obj:`numba.core.dispatcher.Dispatcher` + + A dpex.djit-specific dispatcher. + + The DpjitDispatcher sets the targetdescr string to "dpex" so that Numba's + Dispatcher can lookup the global target_registry with that string and + correctly use the DpexTarget context. + + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`targetdescr ` + - \- + + + + .. rubric:: Members + + .. py:attribute:: targetdescr + + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/exceptions/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/exceptions/index.rst.txt new file mode 100644 index 0000000000..a35111b786 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/exceptions/index.rst.txt @@ -0,0 +1,22 @@ + +:orphan: + +numba_dpex.core.exceptions +========================== + +.. py:module:: numba_dpex.core.exceptions + +.. autoapi-nested-parse:: + + The module defines the custom error classes used in numba_dpex. + + + +Overview +-------- + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/index.rst.txt new file mode 100644 index 0000000000..649e60ac26 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/index.rst.txt @@ -0,0 +1,347 @@ + +numba_dpex.core +=============== + +.. py:module:: numba_dpex.core + + +Subpackages +----------- +.. toctree:: + :titlesonly: + :maxdepth: 3 + + boxing/index.rst + datamodel/index.rst + kernel_interface/index.rst + parfors/index.rst + passes/index.rst + pipelines/index.rst + runtime/index.rst + targets/index.rst + typeconv/index.rst + types/index.rst + typing/index.rst + utils/index.rst + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`Array ` + - An array type for use inside our compiler pipeline. + * - :py:obj:`DpctlSyclQueue ` + - A Numba type to represent a dpctl.SyclQueue PyObject. + * - :py:obj:`DpctlSyclEvent ` + - A Numba type to represent a dpctl.SyclEvent PyObject. + * - :py:obj:`DpnpNdArray ` + - The Numba type to represent an dpnp.ndarray. The type has the same + * - :py:obj:`IntEnumLiteral ` + - A Literal type for IntEnum objects. The type contains the original Python + * - :py:obj:`NdRangeType ` + - Numba-dpex type corresponding to + * - :py:obj:`RangeType ` + - Numba-dpex type corresponding to + * - :py:obj:`USMNdArray ` + - A type class to represent dpctl.tensor.usm_ndarray. + + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`none ` + - \- + * - :py:obj:`uint32 ` + - \- + * - :py:obj:`uint64 ` + - \- + * - :py:obj:`int32 ` + - \- + * - :py:obj:`int64 ` + - \- + * - :py:obj:`float32 ` + - \- + * - :py:obj:`float64 ` + - \- + * - :py:obj:`b1 ` + - \- + * - :py:obj:`i4 ` + - \- + * - :py:obj:`i8 ` + - \- + * - :py:obj:`u4 ` + - \- + * - :py:obj:`u8 ` + - \- + * - :py:obj:`f4 ` + - \- + * - :py:obj:`f8 ` + - \- + * - :py:obj:`float_ ` + - \- + * - :py:obj:`double ` + - \- + * - :py:obj:`usm_ndarray ` + - \- + * - :py:obj:`void ` + - \- + + +Classes +------- + +.. py:class:: Array(dtype, ndim, layout, readonly=False, name=None, aligned=True, addrspace=None) + + Bases: :py:obj:`Array` + + An array type for use inside our compiler pipeline. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy `\ (dtype, ndim, layout, readonly, addrspace) + - \- + * - :py:obj:`is_precise `\ () + - Whether this type is precise, i.e. can be part of a successful + + + .. rubric:: Members + + .. py:method:: copy(dtype=None, ndim=None, layout=None, readonly=None, addrspace=None) + + + .. py:method:: is_precise() + + Whether this type is precise, i.e. can be part of a successful + type inference. Default implementation returns True. + + + + +.. py:class:: DpctlSyclQueue(sycl_queue) + + Bases: :py:obj:`numba.types.Type` + + A Numba type to represent a dpctl.SyclQueue PyObject. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`rand_digit_str `\ (n) + - \- + + + .. rubric:: Members + + .. py:method:: rand_digit_str(n) + + + + +.. py:class:: DpctlSyclEvent + + Bases: :py:obj:`numba.types.Type` + + A Numba type to represent a dpctl.SyclEvent PyObject. + + + + +.. py:class:: DpnpNdArray(ndim, layout='C', dtype=None, usm_type='device', device=None, queue=None, readonly=False, name=None, aligned=True, addrspace=address_space.GLOBAL) + + Bases: :py:obj:`numba_dpex.core.types.usm_ndarray_type.USMNdArray` + + The Numba type to represent an dpnp.ndarray. The type has the same + structure as USMNdArray used to represent dpctl.tensor.usm_ndarray. + + + + +.. py:class:: IntEnumLiteral(value) + + Bases: :py:obj:`numba.core.types.Literal`, :py:obj:`numba.core.types.Integer` + + A Literal type for IntEnum objects. The type contains the original Python + value of the IntEnum class in it. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`can_convert_to `\ (typingctx, other) + - Check whether this type can be converted to the *other*. + + + .. rubric:: Members + + .. py:method:: can_convert_to(typingctx, other) -> bool + + Check whether this type can be converted to the *other*. + If successful, must return a string describing the conversion, e.g. + "exact", "promote", "unsafe", "safe"; otherwise None is returned. + + + + +.. py:class:: NdRangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.NdRange` + + + + +.. py:class:: RangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.Range` + + + + +.. py:class:: USMNdArray(ndim, layout='C', dtype=None, usm_type='device', device=None, queue=None, readonly=False, name=None, aligned=True, addrspace=address_space.GLOBAL) + + Bases: :py:obj:`numba.core.types.npytypes.Array` + + A type class to represent dpctl.tensor.usm_ndarray. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy `\ (dtype, ndim, layout, readonly, addrspace, device, usm_type) + - \- + * - :py:obj:`unify `\ (typingctx, other) + - Unify this with the *other* USMNdArray. + * - :py:obj:`can_convert_to `\ (typingctx, other) + - Convert this USMNdArray to the *other*. + + + .. rubric:: Members + + .. py:method:: copy(dtype=None, ndim=None, layout=None, readonly=None, addrspace=None, device=None, usm_type=None) + + + .. py:method:: unify(typingctx, other) + + Unify this with the *other* USMNdArray. + + + .. py:method:: can_convert_to(typingctx, other) + + Convert this USMNdArray to the *other*. + + + + + +Attributes +---------- +.. py:data:: none + + + +.. py:data:: uint32 + + + +.. py:data:: uint64 + + + +.. py:data:: int32 + + + +.. py:data:: int64 + + + +.. py:data:: float32 + + + +.. py:data:: float64 + + + +.. py:data:: b1 + + + +.. py:data:: i4 + + + +.. py:data:: i8 + + + +.. py:data:: u4 + + + +.. py:data:: u8 + + + +.. py:data:: f4 + + + +.. py:data:: f8 + + + +.. py:data:: float_ + + + +.. py:data:: double + + + +.. py:data:: usm_ndarray + + + +.. py:data:: void + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/itanium_mangler/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/itanium_mangler/index.rst.txt new file mode 100644 index 0000000000..1d618454eb --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/itanium_mangler/index.rst.txt @@ -0,0 +1,201 @@ + +:orphan: + +numba_dpex.core.itanium_mangler +=============================== + +.. py:module:: numba_dpex.core.itanium_mangler + +.. autoapi-nested-parse:: + + Itanium CXX ABI Mangler + + Reference: https://itanium-cxx-abi.github.io/cxx-abi/abi.html + + The basics of the mangling scheme. + + We are hijacking the CXX mangling scheme for our use. We map Python modules + into CXX namespace. A `module1.submodule2.foo` is mapped to + `module1::submodule2::foo`. For parameterized numba types, we treat them as + templated types; for example, `array(int64, 1d, C)` becomes an + `array`. + + All mangled names are prefixed with "_Z". It is followed by the name of the + entity. A name contains one or more identifiers. Each identifier is encoded + as "". If the name is namespaced and, therefore, + has multiple identifiers, the entire name is encoded as "NE". + + For functions, arguments types follow. There are condensed encodings for basic + built-in types; e.g. "i" for int, "f" for float. For other types, the + previously mentioned name encoding should be used. + + For templated types, the template parameters are encoded immediately after the + name. If it is namespaced, it should be within the 'N' 'E' marker. Template + parameters are encoded in "IE", where each parameter is encoded using + the mentioned name encoding scheme. Template parameters can contain literal + values like the '1' in the array type shown earlier. There is special encoding + scheme for them to avoid leading digits. + + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`mangle_type_or_value `\ (typ) + - Mangle type parameter and arbitrary value. + * - :py:obj:`mangle_ext `\ (ident, argtys, \*None, abi_tags) + - Mangle identifier with Numba type objects and abi-tags. + * - :py:obj:`mangle_abi_tag `\ (abi_tag) + - \- + * - :py:obj:`mangle_identifier `\ (ident, template_params, \*None, abi_tags, uid) + - Mangle the identifier with optional template parameters and abi_tags. + * - :py:obj:`mangle_type_c `\ (typ) + - Mangle C type name + * - :py:obj:`mangle_type_or_value_numba `\ (typ) + - Mangle type parameter and arbitrary value. + * - :py:obj:`mangle_templated_ident `\ (identifier, parameters) + - Mangle templated identifier. + * - :py:obj:`mangle_args_c `\ (argtys) + - Mangle sequence of C type names + * - :py:obj:`mangle_args `\ (argtys) + - Mangle sequence of Numba type objects and arbitrary values. + * - :py:obj:`mangle_c `\ (ident, argtys) + - Mangle identifier with C type names + * - :py:obj:`mangle `\ (ident, argtys, \*None, abi_tags, uid) + - Mangle identifier with Numba type objects and abi-tags. + * - :py:obj:`prepend_namespace `\ (mangled, ns) + - Prepend namespace to mangled name. + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`PREFIX ` + - \- + * - :py:obj:`C2CODE ` + - \- + * - :py:obj:`N2C ` + - \- + * - :py:obj:`mangle_type ` + - \- + * - :py:obj:`mangle_value ` + - \- + + + +Functions +--------- +.. py:function:: mangle_type_or_value(typ) + + Mangle type parameter and arbitrary value. + + This function extends Numba's `magle_type_or_value()` to + support numba.types.CPointer type, e.g. an ``int *`` argument will be + mangled to "Pi". + Mangling of extended qualifiers is supported only + for address space qualifiers. In which case, the mangling + follows the rule defined in Section 5.1.5.1 of the ``Itanium ABI + ``_. + For example, an ``int global *`` argument will be mangeled to "PU3AS1i". + + :param typ: Type to mangle + :type typ: numba.types, int, str + + :returns: The mangled name of the type + :rtype: str + + +.. py:function:: mangle_ext(ident, argtys, *, abi_tags=()) + + Mangle identifier with Numba type objects and abi-tags. + + +.. py:function:: mangle_abi_tag(abi_tag: str) -> str + + +.. py:function:: mangle_identifier(ident, template_params='', *, abi_tags=(), uid=None) + + Mangle the identifier with optional template parameters and abi_tags. + + Note: + + This treats '.' as '::' in C++. + + +.. py:function:: mangle_type_c(typ) + + Mangle C type name + + :param typ: C type name + :type typ: str + + +.. py:function:: mangle_type_or_value_numba(typ) + + Mangle type parameter and arbitrary value. + + +.. py:function:: mangle_templated_ident(identifier, parameters) + + Mangle templated identifier. + + +.. py:function:: mangle_args_c(argtys) + + Mangle sequence of C type names + + +.. py:function:: mangle_args(argtys) + + Mangle sequence of Numba type objects and arbitrary values. + + +.. py:function:: mangle_c(ident, argtys) + + Mangle identifier with C type names + + +.. py:function:: mangle(ident, argtys, *, abi_tags=(), uid=None) + + Mangle identifier with Numba type objects and abi-tags. + + +.. py:function:: prepend_namespace(mangled, ns) + + Prepend namespace to mangled name. + + + +Attributes +---------- +.. py:data:: PREFIX + :value: '_Z' + + + +.. py:data:: C2CODE + + + +.. py:data:: N2C + + + +.. py:data:: mangle_type + + + +.. py:data:: mangle_value + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/arg_pack_unpacker/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/arg_pack_unpacker/index.rst.txt new file mode 100644 index 0000000000..9fa0776ca9 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/arg_pack_unpacker/index.rst.txt @@ -0,0 +1,37 @@ + +:orphan: + +numba_dpex.core.kernel_interface.arg_pack_unpacker +================================================== + +.. py:module:: numba_dpex.core.kernel_interface.arg_pack_unpacker + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`Packer ` + - Implements the functionality to unpack a Python object passed as an + + + + +Classes +------- + +.. py:class:: Packer(kernel_name, arg_list, argty_list, queue) + + Implements the functionality to unpack a Python object passed as an + argument to a numba_dpex kernel function into corresponding ctype object. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/arrayobj/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/arrayobj/index.rst.txt new file mode 100644 index 0000000000..dd10f2377a --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/arrayobj/index.rst.txt @@ -0,0 +1,60 @@ + +:orphan: + +numba_dpex.core.kernel_interface.arrayobj +========================================= + +.. py:module:: numba_dpex.core.kernel_interface.arrayobj + +.. autoapi-nested-parse:: + + This package contains implementation of some numpy.np.arrayobj functions without + parent and meminfo fields required, because they don't make sense on device. + These functions intended to be used only in kernel targets like local/private or + usm array view. + + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`populate_array `\ (array, data, shape, strides, itemsize) + - Helper function for populating array structures. + * - :py:obj:`make_view `\ (context, builder, aryty, ary, return_type, data, shapes, strides) + - Build a view over the given array with the given parameters. + + + + +Functions +--------- +.. py:function:: populate_array(array, data, shape, strides, itemsize) + + Helper function for populating array structures. + This avoids forgetting to set fields. + + *shape* and *strides* can be Python tuples or LLVM arrays. + + This is analog of numpy.np.arrayobj.populate_array without parent and + meminfo fields, because they don't make sense on device. This function + intended to be used only in kernel targets. + + +.. py:function:: make_view(context, builder, aryty, ary, return_type, data, shapes, strides) + + Build a view over the given array with the given parameters. + + This is analog of numpy.np.arrayobj.make_view without parent and + meminfo fields, because they don't make sense on device. This function + intended to be used only in kernel targets. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/dispatcher/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/dispatcher/index.rst.txt new file mode 100644 index 0000000000..48a004099e --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/dispatcher/index.rst.txt @@ -0,0 +1,44 @@ + +:orphan: + +numba_dpex.core.kernel_interface.dispatcher +=========================================== + +.. py:module:: numba_dpex.core.kernel_interface.dispatcher + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`JitKernel ` + - Functor to wrap a kernel function and JIT compile and dispatch it to a + + + + +Classes +------- + +.. py:class:: JitKernel(pyfunc, debug_flags=None, compile_flags=None, specialization_sigs=None, enable_cache=True) + + Functor to wrap a kernel function and JIT compile and dispatch it to a + specified SYCL queue. + + A JitKernel is returned by the kernel decorator and wraps an instance of a + device kernel function. A device kernel function is specialized for a + backend may represent a binary object in a lower-level IR. Currently, only + SPIR-V binary format device functions for level-zero and opencl backends + are supported. + + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/func/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/func/index.rst.txt new file mode 100644 index 0000000000..02bc6410fa --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/func/index.rst.txt @@ -0,0 +1,181 @@ + +:orphan: + +numba_dpex.core.kernel_interface.func +===================================== + +.. py:module:: numba_dpex.core.kernel_interface.func + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpexFunction ` + - Class to materialize dpex function + * - :py:obj:`DpexFunctionTemplate ` + - Helper class to compile an unspecialized `numba_dpex.func` + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`compile_func `\ (pyfunc, signature, debug) + - Compiles a specialized `numba_dpex.func` + * - :py:obj:`compile_func_template `\ (pyfunc, debug, enable_cache) + - Converts a `numba_dpex.func` function to an `AbstractTemplate` + + + +Classes +------- + +.. py:class:: DpexFunction(pyfunc, debug=False) + + Bases: :py:obj:`object` + + Class to materialize dpex function + + Helper class to eager compile a specialized `numba_dpex.func` + decorated Python function into a LLVM function with `spir_func` + calling convention. + + A specialized `numba_dpex.func` decorated Python function is one + where the user has specified a signature or a list of signatures + for the function. The function gets compiled as soon as the Python + program is loaded, i.e., eagerly, instead of JIT compilation once + the function is invoked. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`compile `\ (arg_types, return_types) + - The actual compilation function. + + + .. rubric:: Members + + .. py:method:: compile(arg_types, return_types) + + The actual compilation function. + + :param arg_types: Function argument types in a tuple. + :type arg_types: `tuple` + :param return_types: An integer value to specify the return type. + :type return_types: `numba.core.types.scalars.Integer` + + :returns: The compiled result + :rtype: `numba.core.compiler.CompileResult` + + + + +.. py:class:: DpexFunctionTemplate(pyfunc, debug=False, enable_cache=True) + + Bases: :py:obj:`object` + + Helper class to compile an unspecialized `numba_dpex.func` + + A helper class to JIT compile an unspecialized `numba_dpex.func` + decorated Python function into an LLVM function with `spir_func` + calling convention. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`compile `\ (args) + - Compile a `numba_dpex.func` decorated function + + + .. rubric:: Members + + .. py:method:: compile(args) + + Compile a `numba_dpex.func` decorated function + + Compile a `numba_dpex.func` decorated Python function with the + given argument types. Each signature is compiled once by caching + the compiled function inside this object. + + :param args: Function argument types in a tuple. + :type args: `tuple` + + :returns: + + Signature of the + compiled result. + :rtype: `numba.core.typing.templates.Signature` + + + + +Functions +--------- +.. py:function:: compile_func(pyfunc, signature, debug=False) + + Compiles a specialized `numba_dpex.func` + + Compiles a specialized `numba_dpex.func` decorated function to native binary + library function and returns the library wrapped inside a + `numba_dpex.core.kernel_interface.func.DpexFunction` object. + + :param pyfunc: A python function to be compiled. + :type pyfunc: `function` + :param signature: A list of `numba.core.typing.templates.Signature`'s + :type signature: `list` + :param debug: Debug options. Defaults to `False`. + :type debug: `bool`, optional + + :returns: + + A `DpexFunction` + object + :rtype: `numba_dpex.core.kernel_interface.func.DpexFunction` + + +.. py:function:: compile_func_template(pyfunc, debug=False, enable_cache=True) + + Converts a `numba_dpex.func` function to an `AbstractTemplate` + + Converts a `numba_dpex.func` decorated function to a Numba + `AbstractTemplate` and returns the object wrapped inside a + `numba_dpex.core.kernel_interface.func.DpexFunctionTemplate` + object. + + A `DpexFunctionTemplate` object is an abstract representation for + a native function with `spir_func` calling convention that is to be + JIT compiled once the argument types are resolved. + + :param pyfunc: A python function to be compiled. + :type pyfunc: `function` + :param debug: Debug options. Defaults to `False`. + :type debug: `bool`, optional + + :raises AssertionError: Raised if keyword arguments are supplied in + the inner generic function. + + :returns: A `DpexFunctionTemplate` object. + :rtype: `numba_dpex.core.kernel_interface.func.DpexFunctionTemplate` + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/index.rst.txt new file mode 100644 index 0000000000..33d673c7fd --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/index.rst.txt @@ -0,0 +1,12 @@ + +numba_dpex.core.kernel_interface +================================ + +.. py:module:: numba_dpex.core.kernel_interface + +.. autoapi-nested-parse:: + + Defines the interface for kernel compilation using numba-dpex. + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/kernel_base/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/kernel_base/index.rst.txt new file mode 100644 index 0000000000..17acf49f2b --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/kernel_base/index.rst.txt @@ -0,0 +1,63 @@ + +:orphan: + +numba_dpex.core.kernel_interface.kernel_base +============================================ + +.. py:module:: numba_dpex.core.kernel_interface.kernel_base + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`KernelInterface ` + - An interface for compute kernel that was generated either from a + + + + +Classes +------- + +.. py:class:: KernelInterface + + An interface for compute kernel that was generated either from a + Python function object or as a Numba IR FunctionType object. + + :param metaclass: The interface is derived from abc.ABCMeta. + :type metaclass: optional + + :raises NotImplementedError: The interface does not implement any of the + :raises methods and subclasses are required to implement them.: + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`compile `\ (target_ctx, typing_ctx, args, debug, compile_flags) + - :summarylabel:`abc` Abstract method to compile a Kernel instance. + + + .. rubric:: Members + + .. py:method:: compile(target_ctx, typing_ctx, args, debug, compile_flags) + :abstractmethod: + + Abstract method to compile a Kernel instance. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/launcher/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/launcher/index.rst.txt new file mode 100644 index 0000000000..81eb8f2eea --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/launcher/index.rst.txt @@ -0,0 +1,51 @@ + +:orphan: + +numba_dpex.core.kernel_interface.launcher +========================================= + +.. py:module:: numba_dpex.core.kernel_interface.launcher + +.. autoapi-nested-parse:: + + Launcher package to provide the same way of calling kernel as experimental + one. + + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`call_kernel `\ (kernel_fn, index_space, \*kernel_args) + - Syntax sugar for calling kernel the same way as experimental one. + + + + +Functions +--------- +.. py:function:: call_kernel(kernel_fn, index_space, *kernel_args) -> None + + Syntax sugar for calling kernel the same way as experimental one. + It is a temporary glue for the experimental kernel migration. + + :param kernel_fn: A + :type kernel_fn: numba_dpex.experimental.KernelDispatcher + :param numba_dpex.kernel decorated function that is compiled to a: + :param KernelDispatcher by numba_dpex.: + :param index_space: A numba_dpex.Range or numba_dpex.NdRange + :type index_space: Range | NdRange + :param type object that specifies the index space for the kernel.: + :param kernel_args: List of objects that are passed to the numba_dpex.kernel + :param decorated function.: + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/ranges_overloads/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/ranges_overloads/index.rst.txt new file mode 100644 index 0000000000..4dbb25ef25 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/ranges_overloads/index.rst.txt @@ -0,0 +1,24 @@ + +:orphan: + +numba_dpex.core.kernel_interface.ranges_overloads +================================================= + +.. py:module:: numba_dpex.core.kernel_interface.ranges_overloads + + +Overview +-------- + + + + +Attributes +---------- +.. py:data:: DPEX_TARGET_NAME + :value: 'dpex' + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/spirv_kernel/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/spirv_kernel/index.rst.txt new file mode 100644 index 0000000000..8db464f550 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/spirv_kernel/index.rst.txt @@ -0,0 +1,71 @@ + +:orphan: + +numba_dpex.core.kernel_interface.spirv_kernel +============================================= + +.. py:module:: numba_dpex.core.kernel_interface.spirv_kernel + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`SpirvKernel ` + - An interface for compute kernel that was generated either from a + + + + +Classes +------- + +.. py:class:: SpirvKernel(func, func_name) + + Bases: :py:obj:`numba_dpex.core.kernel_interface.kernel_base.KernelInterface` + + An interface for compute kernel that was generated either from a + Python function object or as a Numba IR FunctionType object. + + :param metaclass: The interface is derived from abc.ABCMeta. + :type metaclass: optional + + :raises NotImplementedError: The interface does not implement any of the + :raises methods and subclasses are required to implement them.: + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`compile `\ (target_ctx, typing_ctx, args, debug, compile_flags) + - Compiles a kernel using numba_dpex.core.compiler.Compiler. + + + .. rubric:: Members + + .. py:method:: compile(target_ctx, typing_ctx, args, debug, compile_flags) + + Compiles a kernel using numba_dpex.core.compiler.Compiler. + + :param args: _description_ + :type args: _type_ + :param debug: _description_ + :type debug: _type_ + :param compile_flags: _description_ + :type compile_flags: _type_ + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/kernel_interface/utils/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/utils/index.rst.txt new file mode 100644 index 0000000000..3b531403e3 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_interface/utils/index.rst.txt @@ -0,0 +1,66 @@ + +:orphan: + +numba_dpex.core.kernel_interface.utils +====================================== + +.. py:module:: numba_dpex.core.kernel_interface.utils + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`chk_compute_follows_data_compliance `\ (usm_array_arglist) + - Check if all the usm ndarray's have the same device. + * - :py:obj:`determine_kernel_launch_queue `\ (args, argtypes, kernel_name) + - Determines the queue where the kernel is to be launched. + + + + +Functions +--------- +.. py:function:: chk_compute_follows_data_compliance(usm_array_arglist) + + Check if all the usm ndarray's have the same device. + + Extracts the device filter string from the Numba inferred USMNdArray + type. Check if the devices corresponding to the filter string are + equivalent and return a ``dpctl.SyclDevice`` object corresponding to the + common filter string. + + If an exception occurred in creating a ``dpctl.SyclDevice``, or the + devices are not equivalent then returns None. + + :param usm_array_arglist: A list of usm_ndarray types specified as + :param arguments to the kernel.: + + :returns: A ``dpctl.SyclDevice`` object if all USMNdArray have same device, or + else None is returned. + + +.. py:function:: determine_kernel_launch_queue(args, argtypes, kernel_name) + + Determines the queue where the kernel is to be launched. + + The execution queue is derived following Python Array API's + "compute follows data" programming model. + + :param argtypes: The Numba inferred type for each argument. + :param kernel_name: The name of the kernel function + + :returns: A queue the common queue used to allocate the arrays. If no such + queue exists, then raises an Exception. + + :raises ExecutionQueueInferenceError: If the queue could not be inferred. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/experimental/launcher/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/kernel_launcher/index.rst.txt similarity index 88% rename from dev/_sources/autoapi/numba_dpex/experimental/launcher/index.rst.txt rename to dev/_sources/autoapi/numba_dpex/core/kernel_launcher/index.rst.txt index 35c749053c..49ede32bcd 100644 --- a/dev/_sources/autoapi/numba_dpex/experimental/launcher/index.rst.txt +++ b/dev/_sources/autoapi/numba_dpex/core/kernel_launcher/index.rst.txt @@ -1,10 +1,10 @@ :orphan: -numba_dpex.experimental.launcher -================================ +numba_dpex.core.kernel_launcher +=============================== -.. py:module:: numba_dpex.experimental.launcher +.. py:module:: numba_dpex.core.kernel_launcher .. autoapi-nested-parse:: @@ -22,9 +22,9 @@ Overview :widths: auto :class: summarytable - * - :py:obj:`call_kernel `\ (kernel_fn, index_space, \*kernel_args) + * - :py:obj:`call_kernel `\ (kernel_fn, index_space, \*kernel_args) - Compiles and synchronously executes a kernel function. - * - :py:obj:`call_kernel_async `\ (kernel_fn, index_space, dependent_events, \*kernel_args) + * - :py:obj:`call_kernel_async `\ (kernel_fn, index_space, dependent_events, \*kernel_args) - Compiles and asynchronously executes a kernel function. diff --git a/dev/_sources/autoapi/numba_dpex/core/parfors/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/parfors/index.rst.txt new file mode 100644 index 0000000000..3180afdad0 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/parfors/index.rst.txt @@ -0,0 +1,7 @@ + +numba_dpex.core.parfors +======================= + +.. py:module:: numba_dpex.core.parfors + + diff --git a/dev/_sources/autoapi/numba_dpex/core/parfors/kernel_builder/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/parfors/kernel_builder/index.rst.txt new file mode 100644 index 0000000000..cefbff578d --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/parfors/kernel_builder/index.rst.txt @@ -0,0 +1,79 @@ + +:orphan: + +numba_dpex.core.parfors.kernel_builder +====================================== + +.. py:module:: numba_dpex.core.parfors.kernel_builder + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`ParforKernel ` + - \- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`create_kernel_for_parfor `\ (lowerer, parfor_node, typemap, flags, loop_ranges, has_aliases, races, parfor_outputs) + - Creates a numba_dpex.kernel function for a parfor node. + * - :py:obj:`update_sentinel `\ (kernel_ir, sentinel_name, kernel_body, new_label) + - Searched all the blocks in the IR generated from a kernel template and + + + +Classes +------- + +.. py:class:: ParforKernel(name, kernel, signature, kernel_args, kernel_arg_types, queue: dpctl.SyclQueue) + + + + +Functions +--------- +.. py:function:: create_kernel_for_parfor(lowerer, parfor_node, typemap, flags, loop_ranges, has_aliases, races, parfor_outputs) -> ParforKernel + + Creates a numba_dpex.kernel function for a parfor node. + + There are two parts to this function: + + 1) Code to iterate across the iteration space as defined by + the schedule. + 2) The parfor body that does the work for a single point in + the iteration space. + + Part 1 is created as Python text for simplicity with a sentinel + assignment to mark the point in the IR where the parfor body + should be added. This Python text is 'exec'ed into existence and its + IR retrieved with run_frontend. The IR is scanned for the sentinel + assignment where that basic block is split and the IR for the parfor + body inserted. + + +.. py:function:: update_sentinel(kernel_ir, sentinel_name, kernel_body, new_label) + + Searched all the blocks in the IR generated from a kernel template and + replaces the __sentinel__ instruction with the actual op for the parfor. + + :param kernel_ir: Numba FunctionIR that was generated from a kernel template + :param sentinel_name: The name of the sentinel instruction that is to be + :param replaced.: + :param kernel_body: The function body of the kernel template generated + :param numba_dpex.kernel function: + :param new_label: The new label to be used for the basic block created to store + :param the instructions that replaced the sentinel: + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/parfors/parfor_lowerer/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/parfors/parfor_lowerer/index.rst.txt new file mode 100644 index 0000000000..c905e9ec52 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/parfors/parfor_lowerer/index.rst.txt @@ -0,0 +1,105 @@ + +:orphan: + +numba_dpex.core.parfors.parfor_lowerer +====================================== + +.. py:module:: numba_dpex.core.parfors.parfor_lowerer + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`ParforLowerImpl ` + - Provides a custom lowerer for parfor nodes that generates a SYCL kernel + * - :py:obj:`ParforLowerFactory ` + - A pseudo-factory class that maps a device filter string to a lowering + + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`keep_alive_kernels ` + - \- + + +Classes +------- + +.. py:class:: ParforLowerImpl + + Provides a custom lowerer for parfor nodes that generates a SYCL kernel + for a parfor and submits it to a queue. + + + + +.. py:class:: ParforLowerFactory + + A pseudo-factory class that maps a device filter string to a lowering + function. + + Each Parfor instruction can have an optional "lowerer" attribute. The + lowerer attribute determines how the parfor instruction should be lowered + to LLVM IR. In addition, the lower attribute decides which parfor + instructions can be fused together. + + The factory class maintains a dictionary mapping every device + type (filter string) encountered so far to a lowerer function for that + device type. At this point numba-dpex does not generate device-specific code + and the lowerer used is same for all device types. However, as a different + ParforLowerImpl instance is returned for every parfor instruction that has + a distinct compute-follows-data inferred device it prevents illegal + parfor fusion. + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`device_to_lowerer_map ` + - \- + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get_lowerer `\ (device) + - :summarylabel:`class` \- + + + .. rubric:: Members + + .. py:attribute:: device_to_lowerer_map + + + + .. py:method:: get_lowerer(device) + :classmethod: + + + + + +Attributes +---------- +.. py:data:: keep_alive_kernels + :value: [] + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/parfors/reduction_helper/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/parfors/reduction_helper/index.rst.txt new file mode 100644 index 0000000000..62eaebcb34 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/parfors/reduction_helper/index.rst.txt @@ -0,0 +1,65 @@ + +:orphan: + +numba_dpex.core.parfors.reduction_helper +======================================== + +.. py:module:: numba_dpex.core.parfors.reduction_helper + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`ReductionHelper ` + - The class to define and allocate reduction intermediate variables. + * - :py:obj:`ReductionKernelVariables ` + - The parfor body and the main function body share ir.Var nodes. + + + + +Classes +------- + +.. py:class:: ReductionHelper + + The class to define and allocate reduction intermediate variables. + + + + +.. py:class:: ReductionKernelVariables(lowerer, parfor_node, typemap, parfor_outputs, reductionHelperList) + + The parfor body and the main function body share ir.Var nodes. + We have to do some replacements of Var names in the parfor body + to make them legal parameter names. If we don't copy then the + Vars in the main function also would incorrectly change their name. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy_final_sum_to_host `\ (parfor_kernel) + - \- + + + .. rubric:: Members + + .. py:method:: copy_final_sum_to_host(parfor_kernel) + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/parfors/reduction_kernel_builder/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/parfors/reduction_kernel_builder/index.rst.txt new file mode 100644 index 0000000000..4cc264fd61 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/parfors/reduction_kernel_builder/index.rst.txt @@ -0,0 +1,41 @@ + +:orphan: + +numba_dpex.core.parfors.reduction_kernel_builder +================================================ + +.. py:module:: numba_dpex.core.parfors.reduction_kernel_builder + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`create_reduction_main_kernel_for_parfor `\ (loop_ranges, parfor_node, typemap, flags, has_aliases, reductionKernelVar, parfor_reddict) + - Creates a numba_dpex.kernel function for reduction main kernel. + * - :py:obj:`create_reduction_remainder_kernel_for_parfor `\ (parfor_node, typemap, flags, has_aliases, reductionKernelVar, parfor_reddict, reductionHelperList) + - Creates a numba_dpex.kernel function for a reduction remainder kernel. + + + + +Functions +--------- +.. py:function:: create_reduction_main_kernel_for_parfor(loop_ranges, parfor_node, typemap, flags, has_aliases, reductionKernelVar, parfor_reddict=None) + + Creates a numba_dpex.kernel function for reduction main kernel. + + +.. py:function:: create_reduction_remainder_kernel_for_parfor(parfor_node, typemap, flags, has_aliases, reductionKernelVar, parfor_reddict, reductionHelperList) + + Creates a numba_dpex.kernel function for a reduction remainder kernel. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/passes/dufunc_inliner/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/passes/dufunc_inliner/index.rst.txt new file mode 100644 index 0000000000..08a1087d76 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/passes/dufunc_inliner/index.rst.txt @@ -0,0 +1,32 @@ + +:orphan: + +numba_dpex.core.passes.dufunc_inliner +===================================== + +.. py:module:: numba_dpex.core.passes.dufunc_inliner + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dufunc_inliner `\ (func_ir, calltypes, typemap, typingctx, targetctx) + - \- + + + + +Functions +--------- +.. py:function:: dufunc_inliner(func_ir, calltypes, typemap, typingctx, targetctx) + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/passes/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/passes/index.rst.txt new file mode 100644 index 0000000000..43e5ef55e3 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/passes/index.rst.txt @@ -0,0 +1,119 @@ + +numba_dpex.core.passes +====================== + +.. py:module:: numba_dpex.core.passes + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`ParforLegalizeCFDPass ` + - Base class for function passes + * - :py:obj:`DumpParforDiagnostics ` + - Base class for analysis passes (no modification made to state) + * - :py:obj:`NoPythonBackend ` + - Base class for function passes + + + + +Classes +------- + +.. py:class:: ParforLegalizeCFDPass + + Bases: :py:obj:`numba.core.compiler_machinery.FunctionPass` + + Base class for function passes + + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Legalize CFD of parfor nodes. + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Legalize CFD of parfor nodes. + + + + +.. py:class:: DumpParforDiagnostics + + Bases: :py:obj:`numba.core.compiler_machinery.AnalysisPass` + + Base class for analysis passes (no modification made to state) + + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Runs the pass itself. Must return True/False depending on whether + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Runs the pass itself. Must return True/False depending on whether + statement level modification took place. + + + + +.. py:class:: NoPythonBackend + + Bases: :py:obj:`numba.core.compiler_machinery.FunctionPass` + + Base class for function passes + + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Back-end: Generate LLVM IR from Numba IR, compile to machine code + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Back-end: Generate LLVM IR from Numba IR, compile to machine code + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/passes/parfor_legalize_cfd_pass/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/passes/parfor_legalize_cfd_pass/index.rst.txt new file mode 100644 index 0000000000..df91b25086 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/passes/parfor_legalize_cfd_pass/index.rst.txt @@ -0,0 +1,122 @@ + +:orphan: + +numba_dpex.core.passes.parfor_legalize_cfd_pass +=============================================== + +.. py:module:: numba_dpex.core.passes.parfor_legalize_cfd_pass + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`ParforLegalizeCFDPassImpl ` + - Legalizes the compute-follows-data based device attribute for parfor + * - :py:obj:`ParforLegalizeCFDPass ` + - Base class for function passes + + + + +Classes +------- + +.. py:class:: ParforLegalizeCFDPassImpl(state) + + Legalizes the compute-follows-data based device attribute for parfor + nodes. + + DpnpNdArray array-expressions populate the type of the left-hand-side (LHS) + of each expression as a default DpnpNdArray instance derived from the + __array_ufunc__ method of DpnpNdArray class. The pass fixes the LHS type by + properly applying compute follows data programming model. The pass first + checks if the right-hand-side (RHS) DpnpNdArray arguments are on the same + device, else raising a ExecutionQueueInferenceError. Once the RHS has + been validated, the LHS type is updated. + + The pass also updated the usm_type of the LHS based on a USM type + propagation rule: device > shared > host. Thus, if the usm_type attribute of + the RHS arrays are "device" and "shared" respectively, the LHS array's + usm_type attribute will be "device". + + Once the pass has identified a parfor with DpnpNdArrays and legalized it, + the "lowerer" attribute of the parfor is set to + ``numba_dpex.core.passes.parfor_lowering_pass._lower_parfor_as_kernel`` so + that the parfor node is lowered using Dpex's lowerer. + + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`inputUsmTypeStrToInt ` + - \- + * - :py:obj:`inputUsmTypeIntToStr ` + - \- + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run `\ () + - \- + + + .. rubric:: Members + + .. py:attribute:: inputUsmTypeStrToInt + + + + .. py:attribute:: inputUsmTypeIntToStr + + + + .. py:method:: run() + + + + +.. py:class:: ParforLegalizeCFDPass + + Bases: :py:obj:`numba.core.compiler_machinery.FunctionPass` + + Base class for function passes + + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Legalize CFD of parfor nodes. + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Legalize CFD of parfor nodes. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/passes/passes/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/passes/passes/index.rst.txt new file mode 100644 index 0000000000..01dc8c3021 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/passes/passes/index.rst.txt @@ -0,0 +1,164 @@ + +:orphan: + +numba_dpex.core.passes.passes +============================= + +.. py:module:: numba_dpex.core.passes.passes + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`ConstantSizeStaticLocalMemoryPass ` + - Base class for function passes + * - :py:obj:`NoPythonBackend ` + - Base class for function passes + * - :py:obj:`DumpParforDiagnostics ` + - Base class for analysis passes (no modification made to state) + * - :py:obj:`QualNameDisambiguationLowering ` + - Qualified name disambiguation lowering pass + + + + +Classes +------- + +.. py:class:: ConstantSizeStaticLocalMemoryPass + + Bases: :py:obj:`numba.core.compiler_machinery.FunctionPass` + + Base class for function passes + + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Preprocessing for data-parallel computations. + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Preprocessing for data-parallel computations. + + + + +.. py:class:: NoPythonBackend + + Bases: :py:obj:`numba.core.compiler_machinery.FunctionPass` + + Base class for function passes + + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Back-end: Generate LLVM IR from Numba IR, compile to machine code + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Back-end: Generate LLVM IR from Numba IR, compile to machine code + + + + +.. py:class:: DumpParforDiagnostics + + Bases: :py:obj:`numba.core.compiler_machinery.AnalysisPass` + + Base class for analysis passes (no modification made to state) + + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Runs the pass itself. Must return True/False depending on whether + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Runs the pass itself. Must return True/False depending on whether + statement level modification took place. + + + + +.. py:class:: QualNameDisambiguationLowering + + Bases: :py:obj:`numba.core.typed_passes.NativeLowering` + + Qualified name disambiguation lowering pass + + If there are multiple @func decorated functions exist inside + another @func decorated block, the numba compiler machinery + creates same qualified names for different compiled function. + Therefore, we utilize `unique_name` to resolve the ambiguity. + + :param NativeLowering: Superclass from which this + :type NativeLowering: CompilerPass + :param class has been inherited.: + + :returns: True if `run_pass()` of the superclass is successful. + :rtype: bool + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`run_pass `\ (state) + - Runs the pass itself. Must return True/False depending on whether + + + .. rubric:: Members + + .. py:method:: run_pass(state) + + Runs the pass itself. Must return True/False depending on whether + statement level modification took place. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/pipelines/dpjit_compiler/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/pipelines/dpjit_compiler/index.rst.txt new file mode 100644 index 0000000000..01fd689994 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/pipelines/dpjit_compiler/index.rst.txt @@ -0,0 +1,82 @@ + +:orphan: + +numba_dpex.core.pipelines.dpjit_compiler +======================================== + +.. py:module:: numba_dpex.core.pipelines.dpjit_compiler + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpjitCompiler ` + - Dpex's compiler pipeline to offload parfor nodes into SYCL kernels. + * - :py:obj:`DpjitCompilerMlir ` + - Dpex's compiler pipeline to offload parfor nodes into SYCL kernels. + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get_compiler `\ (use_mlir) + - \- + + + +Classes +------- + +.. py:class:: DpjitCompiler(typingctx, targetctx, library, args, return_type, flags, locals) + + Bases: :py:obj:`numba.core.compiler.CompilerBase` + + Dpex's compiler pipeline to offload parfor nodes into SYCL kernels. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`define_pipelines `\ () + - Child classes override this to customize the pipelines in use. + + + .. rubric:: Members + + .. py:method:: define_pipelines() + + Child classes override this to customize the pipelines in use. + + + + + +.. py:class:: DpjitCompilerMlir(typingctx, targetctx, library, args, return_type, flags, locals) + + Bases: :py:obj:`DpjitCompiler` + + Dpex's compiler pipeline to offload parfor nodes into SYCL kernels. + + + + +Functions +--------- +.. py:function:: get_compiler(use_mlir) + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/pipelines/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/pipelines/index.rst.txt new file mode 100644 index 0000000000..cab60cc824 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/pipelines/index.rst.txt @@ -0,0 +1,7 @@ + +numba_dpex.core.pipelines +========================= + +.. py:module:: numba_dpex.core.pipelines + + diff --git a/dev/_sources/autoapi/numba_dpex/core/pipelines/kernel_compiler/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/pipelines/kernel_compiler/index.rst.txt new file mode 100644 index 0000000000..506bff5738 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/pipelines/kernel_compiler/index.rst.txt @@ -0,0 +1,58 @@ + +:orphan: + +numba_dpex.core.pipelines.kernel_compiler +========================================= + +.. py:module:: numba_dpex.core.pipelines.kernel_compiler + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`KernelCompiler ` + - Dpex's kernel compilation pipeline. + + + + +Classes +------- + +.. py:class:: KernelCompiler(typingctx, targetctx, library, args, return_type, flags, locals) + + Bases: :py:obj:`numba.core.compiler.CompilerBase` + + Dpex's kernel compilation pipeline. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`define_pipelines `\ () + - Child classes override this to customize the pipelines in use. + + + .. rubric:: Members + + .. py:method:: define_pipelines() + + Child classes override this to customize the pipelines in use. + + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/runtime/context/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/runtime/context/index.rst.txt new file mode 100644 index 0000000000..70e45972de --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/runtime/context/index.rst.txt @@ -0,0 +1,286 @@ + +:orphan: + +numba_dpex.core.runtime.context +=============================== + +.. py:module:: numba_dpex.core.runtime.context + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpexRTContext ` + - An object providing access to DPEXRT API in the lowering pass. + + + + +Classes +------- + +.. py:class:: DpexRTContext(context) + + Bases: :py:obj:`object` + + An object providing access to DPEXRT API in the lowering pass. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`meminfo_alloc `\ (builder, size, usm_type, queue_ref) + - Wrapper to call :func:`~context.DpexRTContext.meminfo_alloc_unchecked` + * - :py:obj:`meminfo_fill `\ (builder, meminfo, itemsize, dest_is_float, value_is_float, value, queue_ref) + - Wrapper to call :func:`~context.DpexRTContext.meminfo_fill_unchecked` + * - :py:obj:`meminfo_alloc_unchecked `\ (builder, size, usm_type, queue_ref) + - Allocate a new MemInfo with a data payload of `size` bytes. + * - :py:obj:`meminfo_fill_unchecked `\ (builder, meminfo, itemsize, dest_is_float, value_is_float, value, queue_ref) + - Fills an allocated `MemInfo` with the value specified. + * - :py:obj:`arraystruct_from_python `\ (pyapi, obj, ptr) + - Generates a call to DPEXRT_sycl_usm_ndarray_from_python C function + * - :py:obj:`queuestruct_from_python `\ (pyapi, obj, ptr) + - Calls the c function DPEXRT_sycl_queue_from_python + * - :py:obj:`queuestruct_to_python `\ (pyapi, val) + - Calls the c function DPEXRT_sycl_queue_to_python + * - :py:obj:`eventstruct_from_python `\ (pyapi, obj, ptr) + - Calls the c function DPEXRT_sycl_event_from_python + * - :py:obj:`eventstruct_to_python `\ (pyapi, val) + - Calls the c function DPEXRT_sycl_event_to_python + * - :py:obj:`eventstruct_init `\ (pyapi, event, struct) + - Calls the c function DPEXRT_sycl_event_init + * - :py:obj:`usm_ndarray_to_python_acqref `\ (pyapi, aryty, ary, dtypeptr) + - Boxes a DpnpNdArray native object into a Python dpnp.ndarray. + * - :py:obj:`get_queue_from_filter_string `\ (builder, device) + - Calls DPEXRTQueue_CreateFromFilterString to create a new sycl::queue + * - :py:obj:`submit_range `\ (builder, kref, qref, args, argtys, nargs, range, nrange, depevents, ndepevents) + - Calls DPEXRTQueue_CreateFromFilterString to create a new sycl::queue + * - :py:obj:`submit_ndrange `\ (builder, kref, qref, args, argtys, nargs, grange, lrange, ndims, depevents, ndepevents) + - Calls DPEXRTQueue_CreateFromFilterString to create a new sycl::queue + * - :py:obj:`acquire_meminfo_and_schedule_release `\ (builder, args) + - Inserts LLVM IR to call nrt_acquire_meminfo_and_schedule_release. + * - :py:obj:`build_or_get_kernel `\ (builder, args) + - Inserts LLVM IR to call build_or_get_kernel. + * - :py:obj:`kernel_cache_size `\ (builder) + - Inserts LLVM IR to call kernel_cache_size. + + + .. rubric:: Members + + .. py:method:: meminfo_alloc(builder, size, usm_type, queue_ref) + + Wrapper to call :func:`~context.DpexRTContext.meminfo_alloc_unchecked` + with null checking of the returned value. + + + .. py:method:: meminfo_fill(builder, meminfo, itemsize, dest_is_float, value_is_float, value, queue_ref) + + Wrapper to call :func:`~context.DpexRTContext.meminfo_fill_unchecked` + with null checking of the returned value. + + + .. py:method:: meminfo_alloc_unchecked(builder, size, usm_type, queue_ref) + + Allocate a new MemInfo with a data payload of `size` bytes. + + The result of the call is checked and if it is NULL, i.e. allocation + failed, then a MemoryError is raised. If the allocation succeeded then + a pointer to the MemInfo is returned. + + :param builder: LLVM IR builder. + :type builder: `llvmlite.ir.builder.IRBuilder` + :param size: LLVM uint64 value specifying + the size in bytes for the data payload, i.e. i64 %"arg.allocsize" + :type size: `llvmlite.ir.values.Argument` + :param usm_type: An LLVM Argument object + specifying the type of the usm allocator. The constant value + should match the values in + ``dpctl's`` ``libsyclinterface::DPCTLSyclUSMType`` enum, + i.e. i64 %"arg.usm_type". + :type usm_type: `llvmlite.ir.values.Argument` + :param queue_ref: An LLVM argument value storing + the pointer to the address of the queue object, the object can be + `dpctl.SyclQueue()`, i.e. i8* %"arg.queue". + :type queue_ref: `llvmlite.ir.values.Argument` + + :returns: + + A pointer to the `MemInfo` + is returned from the `DPEXRT_MemInfo_alloc` C function call. + :rtype: ret (`llvmlite.ir.instructions.CallInstr`) + + + .. py:method:: meminfo_fill_unchecked(builder, meminfo, itemsize, dest_is_float, value_is_float, value, queue_ref) + + Fills an allocated `MemInfo` with the value specified. + + The result of the call is checked and if it is `NULL`, i.e. the fill + operation failed, then a `MemoryError` is raised. If the fill operation + is succeeded then a pointer to the `MemInfo` is returned. + + :param builder: LLVM IR builder. + :type builder: `llvmlite.ir.builder.IRBuilder` + :param meminfo: LLVM uint64 value + specifying the size in bytes for the data payload. + :type meminfo: `llvmlite.ir.instructions.LoadInstr` + :param itemsize: An LLVM Constant value + specifying the size of the each data item allocated by the + usm allocator. + :type itemsize: `llvmlite.ir.values.Constant` + :param dest_is_float: An LLVM Constant + value specifying if the destination array type is floating + point. + :type dest_is_float: `llvmlite.ir.values.Constant` + :param value_is_float: An LLVM Constant + value specifying if the input value is a floating point. + :type value_is_float: `llvmlite.ir.values.Constant` + :param value: An LLVM Constant value + specifying if the input value that will be used to fill + the array. + :type value: `llvmlite.ir.values.Constant` + :param queue_ref: An LLVM ExtractValue + instruction object to extract the pointer to the queue from the + DpctlSyclQueue type, i.e. %".74" = extractvalue {i8*, i8*} %".73", 1. + :type queue_ref: `llvmlite.ir.instructions.ExtractValue` + + :returns: + + A pointer to the `MemInfo` + is returned from the `DPEXRT_MemInfo_fill` C function call. + :rtype: ret (`llvmlite.ir.instructions.CallInstr`) + + + .. py:method:: arraystruct_from_python(pyapi, obj, ptr) + + Generates a call to DPEXRT_sycl_usm_ndarray_from_python C function + defined in the _DPREXRT_python Python extension. + + + + .. py:method:: queuestruct_from_python(pyapi, obj, ptr) + + Calls the c function DPEXRT_sycl_queue_from_python + + + .. py:method:: queuestruct_to_python(pyapi, val) + + Calls the c function DPEXRT_sycl_queue_to_python + + + .. py:method:: eventstruct_from_python(pyapi, obj, ptr) + + Calls the c function DPEXRT_sycl_event_from_python + + + .. py:method:: eventstruct_to_python(pyapi, val) + + Calls the c function DPEXRT_sycl_event_to_python + + + .. py:method:: eventstruct_init(pyapi, event, struct) + + Calls the c function DPEXRT_sycl_event_init + + + .. py:method:: usm_ndarray_to_python_acqref(pyapi, aryty, ary, dtypeptr) + + Boxes a DpnpNdArray native object into a Python dpnp.ndarray. + + :param pyapi: _description_ + :type pyapi: _type_ + :param aryty: _description_ + :type aryty: _type_ + :param ary: _description_ + :type ary: _type_ + :param dtypeptr: _description_ + :type dtypeptr: _type_ + + :returns: _description_ + :rtype: _type_ + + + .. py:method:: get_queue_from_filter_string(builder, device) + + Calls DPEXRTQueue_CreateFromFilterString to create a new sycl::queue + from a given filter string. + + :param device: An LLVM ArrayType + storing a const string for a DPC++ filter selector string. + :type device: llvmlite.ir.values.FormattedConstant + + Returns: A DPCTLSyclQueueRef pointer. + + + .. py:method:: submit_range(builder, kref, qref, args, argtys, nargs, range, nrange, depevents, ndepevents) + + Calls DPEXRTQueue_CreateFromFilterString to create a new sycl::queue + from a given filter string. + + Returns: A DPCTLSyclQueueRef pointer. + + + .. py:method:: submit_ndrange(builder, kref, qref, args, argtys, nargs, grange, lrange, ndims, depevents, ndepevents) + + Calls DPEXRTQueue_CreateFromFilterString to create a new sycl::queue + from a given filter string. + + Returns: A LLVM IR call inst. + + + .. py:method:: acquire_meminfo_and_schedule_release(builder: llvmlite.ir.IRBuilder, args) + + Inserts LLVM IR to call nrt_acquire_meminfo_and_schedule_release. + + DPCTLSyclEventRef + DPEXRT_nrt_acquire_meminfo_and_schedule_release( + NRT_api_functions *nrt, + DPCTLSyclQueueRef QRef, + NRT_MemInfo **meminfo_array, + size_t meminfo_array_size, + DPCTLSyclEventRef *depERefs, + size_t nDepERefs, + int *status, + ); + + + + .. py:method:: build_or_get_kernel(builder: llvmlite.ir.IRBuilder, args) + + Inserts LLVM IR to call build_or_get_kernel. + + DPCTLSyclKernelRef + DPEXRT_build_or_get_kernel( + const DPCTLSyclContextRef ctx, + const DPCTLSyclDeviceRef dev, + size_t il_hash, + const char *il, + size_t il_length, + const char *compile_opts, + const char *kernel_name, + ); + + + + .. py:method:: kernel_cache_size(builder: llvmlite.ir.IRBuilder) + + Inserts LLVM IR to call kernel_cache_size. + + size_t DPEXRT_kernel_cache_size(); + + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/runtime/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/runtime/index.rst.txt new file mode 100644 index 0000000000..8c78c28b6b --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/runtime/index.rst.txt @@ -0,0 +1,7 @@ + +numba_dpex.core.runtime +======================= + +.. py:module:: numba_dpex.core.runtime + + diff --git a/dev/_sources/autoapi/numba_dpex/core/targets/dpjit_target/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/targets/dpjit_target/index.rst.txt new file mode 100644 index 0000000000..daf6fb305a --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/targets/dpjit_target/index.rst.txt @@ -0,0 +1,155 @@ + +:orphan: + +numba_dpex.core.targets.dpjit_target +==================================== + +.. py:module:: numba_dpex.core.targets.dpjit_target + +.. autoapi-nested-parse:: + + Defines the target and typing contexts for numba_dpex's dpjit decorator. + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`Dpex ` + - Mark the target as CPU. + * - :py:obj:`DpexTypingContext ` + - Custom typing context to support dpjit compilation. + * - :py:obj:`DpexTargetContext ` + - Changes BaseContext calling convention + + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DPEX_TARGET_NAME ` + - \- + * - :py:obj:`dpex_function_registry ` + - \- + + +Classes +------- + +.. py:class:: Dpex + + Bases: :py:obj:`numba.core.target_extension.CPU` + + Mark the target as CPU. + + + + + +.. py:class:: DpexTypingContext + + Bases: :py:obj:`numba.core.typing.Context` + + Custom typing context to support dpjit compilation. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`load_additional_registries `\ () + - Register dpjit specific functions like dpnp ufuncs. + + + .. rubric:: Members + + .. py:method:: load_additional_registries() + + Register dpjit specific functions like dpnp ufuncs. + + + + +.. py:class:: DpexTargetContext(typingctx, target=DPEX_TARGET_NAME) + + Bases: :py:obj:`numba.core.cpu.CPUContext` + + Changes BaseContext calling convention + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`init `\ () + - For subclasses to add initializer + * - :py:obj:`dpexrt `\ () + - \- + * - :py:obj:`load_additional_registries `\ () + - Load dpjit-specific registries. + * - :py:obj:`get_ufunc_info `\ (ufunc_key) + - Get the ufunc implementation for a given ufunc object. + + + .. rubric:: Members + + .. py:method:: init() + + For subclasses to add initializer + + + .. py:method:: dpexrt() + + + .. py:method:: load_additional_registries() + + Load dpjit-specific registries. + + + .. py:method:: get_ufunc_info(ufunc_key) + + Get the ufunc implementation for a given ufunc object. + + The default implementation in BaseContext always raises a + ``NotImplementedError`` exception. Subclasses may raise ``KeyError`` + to signal that the given ``ufunc_key`` is not available. + + :param ufunc_key: + :type ufunc_key: NumPy ufunc + + :returns: **res** -- A mapping of a NumPy ufunc type signature to a lower-level + implementation. + :rtype: dict[str, callable] + + + + + +Attributes +---------- +.. py:data:: DPEX_TARGET_NAME + :value: 'dpex' + + + +.. py:data:: dpex_function_registry + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/targets/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/targets/index.rst.txt new file mode 100644 index 0000000000..bef3acf755 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/targets/index.rst.txt @@ -0,0 +1,7 @@ + +numba_dpex.core.targets +======================= + +.. py:module:: numba_dpex.core.targets + + diff --git a/dev/_sources/autoapi/numba_dpex/core/typeconv/array_conversion/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/typeconv/array_conversion/index.rst.txt new file mode 100644 index 0000000000..3f42e78e4c --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/typeconv/array_conversion/index.rst.txt @@ -0,0 +1,42 @@ + +:orphan: + +numba_dpex.core.typeconv.array_conversion +========================================= + +.. py:module:: numba_dpex.core.typeconv.array_conversion + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`to_usm_ndarray `\ (suai_attrs, addrspace) + - Converts an array-like object that has the _sycl_usm_array_interface__ + + + + +Functions +--------- +.. py:function:: to_usm_ndarray(suai_attrs, addrspace=address_space.GLOBAL) + + Converts an array-like object that has the _sycl_usm_array_interface__ + attribute to numba_dpex.types.UsmNdArray. + + :param suai_attrs: The extracted SUAI information for an array-like object. + :param addrspace: Address space this array is allocated in. + + Returns: The Numba type for SUAI array. + + :raises NotImplementedError: If the dtype of the passed array is not supported. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/typeconv/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/typeconv/index.rst.txt new file mode 100644 index 0000000000..bc9ca9e15d --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/typeconv/index.rst.txt @@ -0,0 +1,40 @@ + +numba_dpex.core.typeconv +======================== + +.. py:module:: numba_dpex.core.typeconv + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`to_usm_ndarray `\ (suai_attrs, addrspace) + - Converts an array-like object that has the _sycl_usm_array_interface__ + + + + +Functions +--------- +.. py:function:: to_usm_ndarray(suai_attrs, addrspace=address_space.GLOBAL) + + Converts an array-like object that has the _sycl_usm_array_interface__ + attribute to numba_dpex.types.UsmNdArray. + + :param suai_attrs: The extracted SUAI information for an array-like object. + :param addrspace: Address space this array is allocated in. + + Returns: The Numba type for SUAI array. + + :raises NotImplementedError: If the dtype of the passed array is not supported. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/array_type/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/array_type/index.rst.txt new file mode 100644 index 0000000000..f604151498 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/array_type/index.rst.txt @@ -0,0 +1,63 @@ + +:orphan: + +numba_dpex.core.types.array_type +================================ + +.. py:module:: numba_dpex.core.types.array_type + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`Array ` + - An array type for use inside our compiler pipeline. + + + + +Classes +------- + +.. py:class:: Array(dtype, ndim, layout, readonly=False, name=None, aligned=True, addrspace=None) + + Bases: :py:obj:`Array` + + An array type for use inside our compiler pipeline. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy `\ (dtype, ndim, layout, readonly, addrspace) + - \- + * - :py:obj:`is_precise `\ () + - Whether this type is precise, i.e. can be part of a successful + + + .. rubric:: Members + + .. py:method:: copy(dtype=None, ndim=None, layout=None, readonly=None, addrspace=None) + + + .. py:method:: is_precise() + + Whether this type is precise, i.e. can be part of a successful + type inference. Default implementation returns True. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/dpctl_types/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/dpctl_types/index.rst.txt new file mode 100644 index 0000000000..bdfff40964 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/dpctl_types/index.rst.txt @@ -0,0 +1,124 @@ + +:orphan: + +numba_dpex.core.types.dpctl_types +================================= + +.. py:module:: numba_dpex.core.types.dpctl_types + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpctlSyclQueue ` + - A Numba type to represent a dpctl.SyclQueue PyObject. + * - :py:obj:`DpctlSyclEvent ` + - A Numba type to represent a dpctl.SyclEvent PyObject. + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`unbox_sycl_queue `\ (typ, obj, c) + - Convert a SyclQueue object to a native structure. + * - :py:obj:`box_sycl_queue `\ (typ, val, c) + - Boxes a NativeValue representation of DpctlSyclQueue type into a + * - :py:obj:`unbox_sycl_event `\ (typ, obj, c) + - Convert a SyclEvent object to a native structure. + * - :py:obj:`box_sycl_event `\ (typ, val, c) + - Boxes a NativeValue representation of DpctlSyclEvent type into a + + + +Classes +------- + +.. py:class:: DpctlSyclQueue(sycl_queue) + + Bases: :py:obj:`numba.types.Type` + + A Numba type to represent a dpctl.SyclQueue PyObject. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`rand_digit_str `\ (n) + - \- + + + .. rubric:: Members + + .. py:method:: rand_digit_str(n) + + + + +.. py:class:: DpctlSyclEvent + + Bases: :py:obj:`numba.types.Type` + + A Numba type to represent a dpctl.SyclEvent PyObject. + + + + +Functions +--------- +.. py:function:: unbox_sycl_queue(typ, obj, c) + + Convert a SyclQueue object to a native structure. + + +.. py:function:: box_sycl_queue(typ, val, c) + + Boxes a NativeValue representation of DpctlSyclQueue type into a + dpctl.SyclQueue PyObject + + At this point numba-dpex does not support creating a dpctl.SyclQueue inside + a dpjit decorated function. For this reason, boxing is only returns the + original parent object stored in DpctlSyclQueue's data model. + + :param typ: The representation of the DpnpNdArray type. + :param val: A native representation of a Numba DpnpNdArray type object. + :param c: The boxing context. + + Returns: A Pyobject for a dpnp.ndarray boxed from the Numba native value. + + +.. py:function:: unbox_sycl_event(typ, obj, c) + + Convert a SyclEvent object to a native structure. + + +.. py:function:: box_sycl_event(typ, val, c) + + Boxes a NativeValue representation of DpctlSyclEvent type into a + dpctl.SyclEvent PyObject + + At this point numba-dpex does not support creating a dpctl.SyclEvent inside + a dpjit decorated function. For this reason, boxing is only returns the + original parent object stored in DpctlSyclEvent's data model. + + :param typ: The representation of the dpctl.SyclEvent type. + :param val: A native representation of a Numba DpctlSyclEvent type object. + :param c: The boxing context. + + Returns: A Pyobject for a dpctl.SyclEvent boxed from the Numba native value. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/dpnp_ndarray_type/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/dpnp_ndarray_type/index.rst.txt new file mode 100644 index 0000000000..3e2b0060ab --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/dpnp_ndarray_type/index.rst.txt @@ -0,0 +1,39 @@ + +:orphan: + +numba_dpex.core.types.dpnp_ndarray_type +======================================= + +.. py:module:: numba_dpex.core.types.dpnp_ndarray_type + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpnpNdArray ` + - The Numba type to represent an dpnp.ndarray. The type has the same + + + + +Classes +------- + +.. py:class:: DpnpNdArray(ndim, layout='C', dtype=None, usm_type='device', device=None, queue=None, readonly=False, name=None, aligned=True, addrspace=address_space.GLOBAL) + + Bases: :py:obj:`numba_dpex.core.types.usm_ndarray_type.USMNdArray` + + The Numba type to represent an dpnp.ndarray. The type has the same + structure as USMNdArray used to represent dpctl.tensor.usm_ndarray. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/index.rst.txt new file mode 100644 index 0000000000..458ffc1fce --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/index.rst.txt @@ -0,0 +1,336 @@ + +numba_dpex.core.types +===================== + +.. py:module:: numba_dpex.core.types + + +Subpackages +----------- +.. toctree:: + :titlesonly: + :maxdepth: 3 + + kernel_api/index.rst + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`Array ` + - An array type for use inside our compiler pipeline. + * - :py:obj:`DpctlSyclEvent ` + - A Numba type to represent a dpctl.SyclEvent PyObject. + * - :py:obj:`DpctlSyclQueue ` + - A Numba type to represent a dpctl.SyclQueue PyObject. + * - :py:obj:`DpnpNdArray ` + - The Numba type to represent an dpnp.ndarray. The type has the same + * - :py:obj:`IntEnumLiteral ` + - A Literal type for IntEnum objects. The type contains the original Python + * - :py:obj:`NdRangeType ` + - Numba-dpex type corresponding to + * - :py:obj:`RangeType ` + - Numba-dpex type corresponding to + * - :py:obj:`USMNdArray ` + - A type class to represent dpctl.tensor.usm_ndarray. + + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`b1 ` + - \- + * - :py:obj:`double ` + - \- + * - :py:obj:`f4 ` + - \- + * - :py:obj:`f8 ` + - \- + * - :py:obj:`float32 ` + - \- + * - :py:obj:`float64 ` + - \- + * - :py:obj:`float_ ` + - \- + * - :py:obj:`i4 ` + - \- + * - :py:obj:`i8 ` + - \- + * - :py:obj:`int32 ` + - \- + * - :py:obj:`int64 ` + - \- + * - :py:obj:`none ` + - \- + * - :py:obj:`u4 ` + - \- + * - :py:obj:`u8 ` + - \- + * - :py:obj:`uint32 ` + - \- + * - :py:obj:`uint64 ` + - \- + * - :py:obj:`void ` + - \- + * - :py:obj:`usm_ndarray ` + - \- + + +Classes +------- + +.. py:class:: Array(dtype, ndim, layout, readonly=False, name=None, aligned=True, addrspace=None) + + Bases: :py:obj:`Array` + + An array type for use inside our compiler pipeline. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy `\ (dtype, ndim, layout, readonly, addrspace) + - \- + * - :py:obj:`is_precise `\ () + - Whether this type is precise, i.e. can be part of a successful + + + .. rubric:: Members + + .. py:method:: copy(dtype=None, ndim=None, layout=None, readonly=None, addrspace=None) + + + .. py:method:: is_precise() + + Whether this type is precise, i.e. can be part of a successful + type inference. Default implementation returns True. + + + + +.. py:class:: DpctlSyclEvent + + Bases: :py:obj:`numba.types.Type` + + A Numba type to represent a dpctl.SyclEvent PyObject. + + + + +.. py:class:: DpctlSyclQueue(sycl_queue) + + Bases: :py:obj:`numba.types.Type` + + A Numba type to represent a dpctl.SyclQueue PyObject. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`rand_digit_str `\ (n) + - \- + + + .. rubric:: Members + + .. py:method:: rand_digit_str(n) + + + + +.. py:class:: DpnpNdArray(ndim, layout='C', dtype=None, usm_type='device', device=None, queue=None, readonly=False, name=None, aligned=True, addrspace=address_space.GLOBAL) + + Bases: :py:obj:`numba_dpex.core.types.usm_ndarray_type.USMNdArray` + + The Numba type to represent an dpnp.ndarray. The type has the same + structure as USMNdArray used to represent dpctl.tensor.usm_ndarray. + + + + +.. py:class:: IntEnumLiteral(value) + + Bases: :py:obj:`numba.core.types.Literal`, :py:obj:`numba.core.types.Integer` + + A Literal type for IntEnum objects. The type contains the original Python + value of the IntEnum class in it. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`can_convert_to `\ (typingctx, other) + - Check whether this type can be converted to the *other*. + + + .. rubric:: Members + + .. py:method:: can_convert_to(typingctx, other) -> bool + + Check whether this type can be converted to the *other*. + If successful, must return a string describing the conversion, e.g. + "exact", "promote", "unsafe", "safe"; otherwise None is returned. + + + + +.. py:class:: NdRangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.NdRange` + + + + +.. py:class:: RangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.Range` + + + + +.. py:class:: USMNdArray(ndim, layout='C', dtype=None, usm_type='device', device=None, queue=None, readonly=False, name=None, aligned=True, addrspace=address_space.GLOBAL) + + Bases: :py:obj:`numba.core.types.npytypes.Array` + + A type class to represent dpctl.tensor.usm_ndarray. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy `\ (dtype, ndim, layout, readonly, addrspace, device, usm_type) + - \- + * - :py:obj:`unify `\ (typingctx, other) + - Unify this with the *other* USMNdArray. + * - :py:obj:`can_convert_to `\ (typingctx, other) + - Convert this USMNdArray to the *other*. + + + .. rubric:: Members + + .. py:method:: copy(dtype=None, ndim=None, layout=None, readonly=None, addrspace=None, device=None, usm_type=None) + + + .. py:method:: unify(typingctx, other) + + Unify this with the *other* USMNdArray. + + + .. py:method:: can_convert_to(typingctx, other) + + Convert this USMNdArray to the *other*. + + + + + +Attributes +---------- +.. py:data:: b1 + + + +.. py:data:: double + + + +.. py:data:: f4 + + + +.. py:data:: f8 + + + +.. py:data:: float32 + + + +.. py:data:: float64 + + + +.. py:data:: float_ + + + +.. py:data:: i4 + + + +.. py:data:: i8 + + + +.. py:data:: int32 + + + +.. py:data:: int64 + + + +.. py:data:: none + + + +.. py:data:: u4 + + + +.. py:data:: u8 + + + +.. py:data:: uint32 + + + +.. py:data:: uint64 + + + +.. py:data:: void + + + +.. py:data:: usm_ndarray + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/atomic_ref/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/atomic_ref/index.rst.txt new file mode 100644 index 0000000000..338e15136b --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/atomic_ref/index.rst.txt @@ -0,0 +1,65 @@ + +:orphan: + +numba_dpex.core.types.kernel_api.atomic_ref +=========================================== + +.. py:module:: numba_dpex.core.types.kernel_api.atomic_ref + +.. autoapi-nested-parse:: + + Collection of numba-dpex typing classes for kernel_api Python classes. + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`AtomicRefType ` + - numba-dpex internal type to represent a Python object of + + + + +Classes +------- + +.. py:class:: AtomicRefType(dtype: numba.core.types.Type, memory_order: int, memory_scope: int, address_space: int) + + Bases: :py:obj:`numba.core.types.Type` + + numba-dpex internal type to represent a Python object of + :class:`numba_dpex.kernel_api.AtomicRef`. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`cast_python_value `\ (args) + - :summarylabel:`abc` The helper function is not overloaded and using it on the + + + .. rubric:: Members + + .. py:method:: cast_python_value(args) + :abstractmethod: + + The helper function is not overloaded and using it on the + AtomicRefType throws a NotImplementedError. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index.rst.txt new file mode 100644 index 0000000000..a8e97da676 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index.rst.txt @@ -0,0 +1,7 @@ + +numba_dpex.core.types.kernel_api +================================ + +.. py:module:: numba_dpex.core.types.kernel_api + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index_space_ids/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index_space_ids/index.rst.txt new file mode 100644 index 0000000000..87ac26b568 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/index_space_ids/index.rst.txt @@ -0,0 +1,119 @@ + +:orphan: + +numba_dpex.core.types.kernel_api.index_space_ids +================================================ + +.. py:module:: numba_dpex.core.types.kernel_api.index_space_ids + +.. autoapi-nested-parse:: + + Defines numba types for Item and NdItem classes + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`GroupType ` + - Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.Group` + * - :py:obj:`ItemType ` + - Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.Item` + * - :py:obj:`NdItemType ` + - Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.NdItem` + + + + +Classes +------- + +.. py:class:: GroupType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.Group` + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`cast_python_value `\ (args) + - :summarylabel:`abc` \- + + + .. rubric:: Members + + .. py:method:: cast_python_value(args) + :abstractmethod: + + + + +.. py:class:: ItemType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.Item` + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`cast_python_value `\ (args) + - :summarylabel:`abc` \- + + + .. rubric:: Members + + .. py:method:: cast_python_value(args) + :abstractmethod: + + + + +.. py:class:: NdItemType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.NdItem` + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`cast_python_value `\ (args) + - :summarylabel:`abc` \- + + + .. rubric:: Members + + .. py:method:: cast_python_value(args) + :abstractmethod: + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/literal_intenum/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/literal_intenum/index.rst.txt new file mode 100644 index 0000000000..10cd950a19 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/literal_intenum/index.rst.txt @@ -0,0 +1,82 @@ + +:orphan: + +numba_dpex.core.types.kernel_api.literal_intenum +================================================ + +.. py:module:: numba_dpex.core.types.kernel_api.literal_intenum + +.. autoapi-nested-parse:: + + Definition of a new Literal type in numba-dpex that allows treating IntEnum + members as integer literals inside a JIT compiled function. + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`IntEnumLiteral ` + - A Literal type for IntEnum objects. The type contains the original Python + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`box_literal_integer `\ (typ, val, ctx) + - Defines how a Numba representation for an IntEnumLiteral object should + + + +Classes +------- + +.. py:class:: IntEnumLiteral(value) + + Bases: :py:obj:`numba.core.types.Literal`, :py:obj:`numba.core.types.Integer` + + A Literal type for IntEnum objects. The type contains the original Python + value of the IntEnum class in it. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`can_convert_to `\ (typingctx, other) + - Check whether this type can be converted to the *other*. + + + .. rubric:: Members + + .. py:method:: can_convert_to(typingctx, other) -> bool + + Check whether this type can be converted to the *other*. + If successful, must return a string describing the conversion, e.g. + "exact", "promote", "unsafe", "safe"; otherwise None is returned. + + + + +Functions +--------- +.. py:function:: box_literal_integer(typ, val, ctx) + + Defines how a Numba representation for an IntEnumLiteral object should + be converted to a PyObject* object and returned back to Python. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/local_accessor/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/local_accessor/index.rst.txt new file mode 100644 index 0000000000..e93f47360e --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/local_accessor/index.rst.txt @@ -0,0 +1,99 @@ + +:orphan: + +numba_dpex.core.types.kernel_api.local_accessor +=============================================== + +.. py:module:: numba_dpex.core.types.kernel_api.local_accessor + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpctlMDLocalAccessorType ` + - numba-dpex internal type to represent a dpctl SyclInterface type + * - :py:obj:`LocalAccessorType ` + - numba-dpex internal type to represent a Python object of + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`unbox_local_accessor `\ (typ, obj, c) + - Unboxes a Python LocalAccessor PyObject* into a numba-dpex internal + + + +Classes +------- + +.. py:class:: DpctlMDLocalAccessorType + + Bases: :py:obj:`numba.core.types.Type` + + numba-dpex internal type to represent a dpctl SyclInterface type + `MDLocalAccessorTy`. + + + + +.. py:class:: LocalAccessorType(ndim, dtype) + + Bases: :py:obj:`numba_dpex.core.types.USMNdArray` + + numba-dpex internal type to represent a Python object of + :class:`numba_dpex.experimental.kernel_iface.LocalAccessor`. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`cast_python_value `\ (args) + - :summarylabel:`abc` The helper function is not overloaded and using it on the + + + .. rubric:: Members + + .. py:method:: cast_python_value(args) + :abstractmethod: + + The helper function is not overloaded and using it on the + LocalAccessorType throws a NotImplementedError. + + + + +Functions +--------- +.. py:function:: unbox_local_accessor(typ, obj, c) + + Unboxes a Python LocalAccessor PyObject* into a numba-dpex internal + representation. + + A LocalAccessor object is represented internally in numba-dpex with the + same data model as a numpy.ndarray. It is done as a LocalAccessor object + serves only as a placeholder type when passed to ``call_kernel`` and the + data buffer should never be accessed inside a host-side compiled function + such as ``call_kernel``. + + When a LocalAccessor object is passed as an argument to a kernel function + it uses the USMArrayDeviceModel. Doing so allows numba-dpex to correctly + generate the kernel signature passing in a pointer in the local address + space. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/ranges/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/ranges/index.rst.txt new file mode 100644 index 0000000000..809bbfe178 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/kernel_api/ranges/index.rst.txt @@ -0,0 +1,51 @@ + +:orphan: + +numba_dpex.core.types.kernel_api.ranges +======================================= + +.. py:module:: numba_dpex.core.types.kernel_api.ranges + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`RangeType ` + - Numba-dpex type corresponding to + * - :py:obj:`NdRangeType ` + - Numba-dpex type corresponding to + + + + +Classes +------- + +.. py:class:: RangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.Range` + + + + +.. py:class:: NdRangeType(ndim: int) + + Bases: :py:obj:`numba.core.types.Type` + + Numba-dpex type corresponding to + :class:`numba_dpex.kernel_api.ranges.NdRange` + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/numba_types_short_names/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/numba_types_short_names/index.rst.txt new file mode 100644 index 0000000000..23588157be --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/numba_types_short_names/index.rst.txt @@ -0,0 +1,87 @@ + +:orphan: + +numba_dpex.core.types.numba_types_short_names +============================================= + +.. py:module:: numba_dpex.core.types.numba_types_short_names + + +Overview +-------- + + + + +Attributes +---------- +.. py:data:: none + + + +.. py:data:: uint32 + + + +.. py:data:: uint64 + + + +.. py:data:: int32 + + + +.. py:data:: int64 + + + +.. py:data:: float32 + + + +.. py:data:: float64 + + + +.. py:data:: b1 + + + +.. py:data:: i4 + + + +.. py:data:: i8 + + + +.. py:data:: u4 + + + +.. py:data:: u8 + + + +.. py:data:: f4 + + + +.. py:data:: f8 + + + +.. py:data:: float_ + + + +.. py:data:: double + + + +.. py:data:: void + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/types/usm_ndarray_type/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/types/usm_ndarray_type/index.rst.txt new file mode 100644 index 0000000000..acaccf9e99 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/types/usm_ndarray_type/index.rst.txt @@ -0,0 +1,74 @@ + +:orphan: + +numba_dpex.core.types.usm_ndarray_type +====================================== + +.. py:module:: numba_dpex.core.types.usm_ndarray_type + +.. autoapi-nested-parse:: + + A type class to represent dpctl.tensor.usm_ndarray type in Numba + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`USMNdArray ` + - A type class to represent dpctl.tensor.usm_ndarray. + + + + +Classes +------- + +.. py:class:: USMNdArray(ndim, layout='C', dtype=None, usm_type='device', device=None, queue=None, readonly=False, name=None, aligned=True, addrspace=address_space.GLOBAL) + + Bases: :py:obj:`numba.core.types.npytypes.Array` + + A type class to represent dpctl.tensor.usm_ndarray. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`copy `\ (dtype, ndim, layout, readonly, addrspace, device, usm_type) + - \- + * - :py:obj:`unify `\ (typingctx, other) + - Unify this with the *other* USMNdArray. + * - :py:obj:`can_convert_to `\ (typingctx, other) + - Convert this USMNdArray to the *other*. + + + .. rubric:: Members + + .. py:method:: copy(dtype=None, ndim=None, layout=None, readonly=None, addrspace=None, device=None, usm_type=None) + + + .. py:method:: unify(typingctx, other) + + Unify this with the *other* USMNdArray. + + + .. py:method:: can_convert_to(typingctx, other) + + Convert this USMNdArray to the *other*. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/typing/dpnpdecl/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/typing/dpnpdecl/index.rst.txt new file mode 100644 index 0000000000..1562fc146f --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/typing/dpnpdecl/index.rst.txt @@ -0,0 +1,164 @@ + +:orphan: + +numba_dpex.core.typing.dpnpdecl +=============================== + +.. py:module:: numba_dpex.core.typing.dpnpdecl + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`DpnpRulesArrayOperator ` + - Defines method ``generic(self, args, kws)`` which compute a possible + * - :py:obj:`DpnpRulesInplaceArrayOperator ` + - Defines method ``generic(self, args, kws)`` which compute a possible + * - :py:obj:`DpnpRulesUnaryArrayOperator ` + - Defines method ``generic(self, args, kws)`` which compute a possible + + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`registry ` + - \- + * - :py:obj:`infer ` + - \- + * - :py:obj:`infer_global ` + - \- + * - :py:obj:`infer_getattr ` + - \- + * - :py:obj:`all_ufuncs ` + - \- + * - :py:obj:`supported_ufuncs ` + - \- + + +Classes +------- + +.. py:class:: DpnpRulesArrayOperator(context) + + Bases: :py:obj:`numba.core.typing.npydecl.NumpyRulesArrayOperator` + + Defines method ``generic(self, args, kws)`` which compute a possible + signature base on input types. The signature does not have to match the + input types. It is compared against the input types afterwards. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`install_operations `\ () + - :summarylabel:`class` \- + + + .. rubric:: Members + + .. py:method:: install_operations() + :classmethod: + + + + +.. py:class:: DpnpRulesInplaceArrayOperator(context) + + Bases: :py:obj:`numba.core.typing.npydecl.NumpyRulesInplaceArrayOperator` + + Defines method ``generic(self, args, kws)`` which compute a possible + signature base on input types. The signature does not have to match the + input types. It is compared against the input types afterwards. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`install_operations `\ () + - :summarylabel:`class` \- + + + .. rubric:: Members + + .. py:method:: install_operations() + :classmethod: + + + + +.. py:class:: DpnpRulesUnaryArrayOperator(context) + + Bases: :py:obj:`numba.core.typing.npydecl.NumpyRulesUnaryArrayOperator` + + Defines method ``generic(self, args, kws)`` which compute a possible + signature base on input types. The signature does not have to match the + input types. It is compared against the input types afterwards. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`install_operations `\ () + - :summarylabel:`class` \- + + + .. rubric:: Members + + .. py:method:: install_operations() + :classmethod: + + + + + +Attributes +---------- +.. py:data:: registry + + + +.. py:data:: infer + + + +.. py:data:: infer_global + + + +.. py:data:: infer_getattr + + + +.. py:data:: all_ufuncs + + + +.. py:data:: supported_ufuncs + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/typing/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/typing/index.rst.txt new file mode 100644 index 0000000000..f312ed4f1d --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/typing/index.rst.txt @@ -0,0 +1,7 @@ + +numba_dpex.core.typing +====================== + +.. py:module:: numba_dpex.core.typing + + diff --git a/dev/_sources/autoapi/numba_dpex/core/typing/typeof/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/typing/typeof/index.rst.txt new file mode 100644 index 0000000000..7cef2d2d73 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/typing/typeof/index.rst.txt @@ -0,0 +1,112 @@ + +:orphan: + +numba_dpex.core.typing.typeof +============================= + +.. py:module:: numba_dpex.core.typing.typeof + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`typeof_usm_ndarray `\ (val, c) + - Registers the type inference implementation function for + * - :py:obj:`typeof_dpnp_ndarray `\ (val, c) + - Registers the type inference implementation function for dpnp.ndarray. + * - :py:obj:`typeof_dpctl_sycl_queue `\ (val, c) + - Registers the type inference implementation function for a + * - :py:obj:`typeof_dpctl_sycl_event `\ (val, c) + - Registers the type inference implementation function for a + * - :py:obj:`typeof_range `\ (val, c) + - Registers the type inference implementation function for a + * - :py:obj:`typeof_ndrange `\ (val, c) + - Registers the type inference implementation function for a + + + + +Functions +--------- +.. py:function:: typeof_usm_ndarray(val, c) + + Registers the type inference implementation function for + dpctl.tensor.usm_ndarray + + :param val: A Python object that should be an instance of a + :param dpctl.tensor.usm_ndarray: + :param c: Unused argument used to be consistent with Numba API. + + :raises ValueError: If an unsupported dtype encountered or val has + :raises no ``usm_type`` or sycl_device attribute.: + + Returns: The Numba type corresponding to dpctl.tensor.usm_ndarray + + +.. py:function:: typeof_dpnp_ndarray(val, c) + + Registers the type inference implementation function for dpnp.ndarray. + + :param val: A Python object that should be an instance of a + :param dpnp.ndarray: + :param c: Unused argument used to be consistent with Numba API. + + :raises ValueError: If an unsupported dtype encountered or val has + :raises no ``usm_type`` or sycl_device attribute.: + + Returns: The Numba type corresponding to dpnp.ndarray + + +.. py:function:: typeof_dpctl_sycl_queue(val, c) + + Registers the type inference implementation function for a + dpctl.SyclQueue PyObject. + + :param val: An instance of dpctl.SyclQueue. + :param c: Unused argument used to be consistent with Numba API. + + Returns: A numba_dpex.core.types.dpctl_types.DpctlSyclQueue instance. + + +.. py:function:: typeof_dpctl_sycl_event(val, c) + + Registers the type inference implementation function for a + dpctl.SyclEvent PyObject. + + :param val: An instance of dpctl.SyclEvent. + :param c: Unused argument used to be consistent with Numba API. + + Returns: A numba_dpex.core.types.dpctl_types.DpctlSyclEvent instance. + + +.. py:function:: typeof_range(val, c) + + Registers the type inference implementation function for a + numba_dpex.Range PyObject. + + :param val: An instance of numba_dpex.Range. + :param c: Unused argument used to be consistent with Numba API. + + Returns: A numba_dpex.core.types.range_types.RangeType instance. + + +.. py:function:: typeof_ndrange(val, c) + + Registers the type inference implementation function for a + numba_dpex.NdRange PyObject. + + :param val: An instance of numba_dpex.Range. + :param c: Unused argument used to be consistent with Numba API. + + Returns: A numba_dpex.core.types.range_types.RangeType instance. + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/caching_utils/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/caching_utils/index.rst.txt new file mode 100644 index 0000000000..71591b264a --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/caching_utils/index.rst.txt @@ -0,0 +1,62 @@ + +:orphan: + +numba_dpex.core.utils.caching_utils +=================================== + +.. py:module:: numba_dpex.core.utils.caching_utils + + +Overview +-------- + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`build_key `\ (\*args) + - Constructs key from variable list of args + * - :py:obj:`create_func_hash `\ (pyfunc) + - Creates a tuple of sha256 hashes out of code and + * - :py:obj:`strip_usm_metadata `\ (argtypes) + - Convert the USMNdArray to an abridged type that disregards the + + + + +Functions +--------- +.. py:function:: build_key(*args) + + Constructs key from variable list of args + + :param \*args: List of components to construct key + + :returns: Tuple of args + + +.. py:function:: create_func_hash(pyfunc) + + Creates a tuple of sha256 hashes out of code and + variable bytes extracted from the compiled funtion. + + :param pyfunc: Python function object + + :returns: Tuple of hashes of code and variable bytes + + +.. py:function:: strip_usm_metadata(argtypes) + + Convert the USMNdArray to an abridged type that disregards the + usm_type, device, queue, address space attributes. + + :param argtypes: List of types + + :returns: Tuple of types after removing USM metadata from USMNdArray type + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/index.rst.txt new file mode 100644 index 0000000000..b0a0fbcde5 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/index.rst.txt @@ -0,0 +1,105 @@ + +numba_dpex.core.utils +===================== + +.. py:module:: numba_dpex.core.utils + + +Subpackages +----------- +.. toctree:: + :titlesonly: + :maxdepth: 3 + + kernel_templates/index.rst + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`SyclUSMArrayInterface ` + - Stores as attributes the information extracted from a + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`build_key `\ (\*args) + - Constructs key from variable list of args + * - :py:obj:`create_func_hash `\ (pyfunc) + - Creates a tuple of sha256 hashes out of code and + * - :py:obj:`strip_usm_metadata `\ (argtypes) + - Convert the USMNdArray to an abridged type that disregards the + * - :py:obj:`get_info_from_suai `\ (obj) + - Extracts the metadata of an object of type UsmNdArray using the objects + + + +Classes +------- + +.. py:class:: SyclUSMArrayInterface(data, writable, size, shape, dimensions, itemsize, strides, dtype, usm_type, device, queue) + + Stores as attributes the information extracted from a + __sycl_usm_array_interface__ dictionary as defined by dpctl.memory.Memory* + classes. + + + + +Functions +--------- +.. py:function:: build_key(*args) + + Constructs key from variable list of args + + :param \*args: List of components to construct key + + :returns: Tuple of args + + +.. py:function:: create_func_hash(pyfunc) + + Creates a tuple of sha256 hashes out of code and + variable bytes extracted from the compiled funtion. + + :param pyfunc: Python function object + + :returns: Tuple of hashes of code and variable bytes + + +.. py:function:: strip_usm_metadata(argtypes) + + Convert the USMNdArray to an abridged type that disregards the + usm_type, device, queue, address space attributes. + + :param argtypes: List of types + + :returns: Tuple of types after removing USM metadata from USMNdArray type + + +.. py:function:: get_info_from_suai(obj) + + Extracts the metadata of an object of type UsmNdArray using the objects + __sycl_usm_array_interface__ (SUAI) attribute. + + The ``dpctl.memory.as_usm_memory`` function converts the array-like + object into a dpctl.memory.USMMemory object. Using the ``as_usm_memory`` + is an implicit way to verify if the array-like object is a legal + SYCL USM memory back Python object that can be passed to a dpex kernel. + + :param obj: array-like object with a SUAI attribute. + + :returns: A SyclUSMArrayInterface object + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/kernel_flattened_args_builder/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_flattened_args_builder/index.rst.txt new file mode 100644 index 0000000000..5fc4e7a599 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_flattened_args_builder/index.rst.txt @@ -0,0 +1,145 @@ + +:orphan: + +numba_dpex.core.utils.kernel_flattened_args_builder +=================================================== + +.. py:module:: numba_dpex.core.utils.kernel_flattened_args_builder + +.. autoapi-nested-parse:: + + Provides helpers to populate the list of kernel arguments that will + be passed to a DPCTLQueue_Submit function call by a KernelLaunchIRBuilder + object. + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`KernelArg ` + - Stores the llvm IR value and the dpctl typeid for a kernel argument. + * - :py:obj:`KernelFlattenedArgsBuilder ` + - Helper to generate the flattened list of kernel arguments to be + + + + +Classes +------- + +.. py:class:: KernelArg + + Bases: :py:obj:`NamedTuple` + + Stores the llvm IR value and the dpctl typeid for a kernel argument. + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`llvm_val ` + - \- + * - :py:obj:`typeid ` + - \- + + + + .. rubric:: Members + + .. py:attribute:: llvm_val + :type: llvmlite.ir.Instruction + + + + .. py:attribute:: typeid + :type: int + + + + + +.. py:class:: KernelFlattenedArgsBuilder(context: numba.core.cpu.CPUContext, builder: llvmlite.ir.IRBuilder, kernel_dmm) + + Helper to generate the flattened list of kernel arguments to be + passed to a DPCTLQueue_Submit function. + + **Note** Two separate data models are used when building a flattened + kernel argument for the following reason: + + Different numba-dpex targets can use different data models for the same + data type that may have different number of attributes and a different + type for each attribute. + + In the case the DpnpNdArray type, two separate data models are used for + the CPUTarget and for the SPIRVTarget. The SPIRVTarget does not have the + ``parent``, ``meminfo`` and ``sycl_queue`` attributes that are present + in the data model used by the CPUTarget. The SPIRVTarget's data model + for DpnpNdArray also requires an explicit address space qualifier for + the ``data`` attribute. + + When generating the LLVM IR for the host-side control code for executing + a SPIR-V kernel, the kernel arguments are represented using the + CPUTarget's data model for each argument's type. However, the actual + kernel function generated as a SPIR-V binary by the SPIRVTarget uses its + own data model manager to build the flattened kernel function argument + list. For this reason, when building the flattened argument list for a + kernel launch call the host data model is used to extract the + required attributes and then the kernel data model is used to get the + correct type for the attribute. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`add_argument `\ (arg_type, arg_packed_llvm_val) + - Add flattened representation of a kernel argument. + * - :py:obj:`get_kernel_arg_list `\ () + - Returns a list of KernelArg objects representing a flattened kernel + * - :py:obj:`print_kernel_arg_list `\ () + - Prints out the kernel argument list in a human readable format. + + + .. rubric:: Members + + .. py:method:: add_argument(arg_type, arg_packed_llvm_val) + + Add flattened representation of a kernel argument. + + + .. py:method:: get_kernel_arg_list() -> list[KernelArg] + + Returns a list of KernelArg objects representing a flattened kernel + argument. + + :returns: List of flattened KernelArg objects + :rtype: list[KernelArg] + + + .. py:method:: print_kernel_arg_list() -> None + + Prints out the kernel argument list in a human readable format. + + :param args_list: List of kernel arguments to be printed + :type args_list: list[KernelArg] + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/kernel_launcher/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_launcher/index.rst.txt new file mode 100644 index 0000000000..fd4bf94958 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_launcher/index.rst.txt @@ -0,0 +1,254 @@ + +:orphan: + +numba_dpex.core.utils.kernel_launcher +===================================== + +.. py:module:: numba_dpex.core.utils.kernel_launcher + +.. autoapi-nested-parse:: + + Module that contains numba style wrapper around sycl kernel submit. + + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`SPIRVKernelModule ` + - Represents SPIRV binary code and function name in this binary + * - :py:obj:`KernelLaunchIRBuilder ` + - Helper class to build the LLVM IR for the submission of a kernel. + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get_queue_from_llvm_values `\ (ctx, builder, ty_kernel_args, ll_kernel_args) + - Get the sycl queue from the first USMNdArray argument. Prior passes + + +.. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`MAX_SIZE_OF_SYCL_RANGE ` + - \- + * - :py:obj:`OPEN_CL_OPT_DISABLE_FLAG ` + - \- + * - :py:obj:`L0_OPT_DISABLE_FLAG ` + - \- + + +Classes +------- + +.. py:class:: SPIRVKernelModule + + Bases: :py:obj:`NamedTuple` + + Represents SPIRV binary code and function name in this binary + + + .. rubric:: Overview + + .. list-table:: Attributes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`kernel_name ` + - \- + * - :py:obj:`kernel_bitcode ` + - \- + + + + .. rubric:: Members + + .. py:attribute:: kernel_name + :type: str + + + + .. py:attribute:: kernel_bitcode + :type: bytes + + + + + +.. py:class:: KernelLaunchIRBuilder(context: numba.core.cpu.CPUContext, builder: llvmlite.ir.builder.IRBuilder, kernel_dmm: numba.core.datamodel.DataModelManager) + + Helper class to build the LLVM IR for the submission of a kernel. + + The class generates LLVM IR inside the current LLVM module that is needed + for submitting kernels. The LLVM Values that + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dpexrt `\ () + - Dpex runtime context. + * - :py:obj:`get_queue `\ (exec_queue) + - Allocates memory on the stack to store a DPCTLSyclQueueRef. + * - :py:obj:`set_kernel `\ (sycl_kernel_ref) + - Sets kernel to the argument list. + * - :py:obj:`set_kernel_from_spirv `\ (kernel_module, debug) + - Sets kernel to the argument list from the SPIRV bytecode. + * - :py:obj:`set_queue `\ (sycl_queue_ref) + - Sets queue to the argument list. + * - :py:obj:`set_queue_from_arguments `\ () + - Sets the sycl queue from the first USMNdArray argument provided + * - :py:obj:`set_range `\ (global_range, local_range) + - Sets global and local range if provided to the argument list. + * - :py:obj:`set_range_from_indexer `\ (ty_indexer_arg, ll_index_arg) + - Returns two lists of LLVM IR Values that hold the unboxed extents of + * - :py:obj:`set_arguments `\ (ty_kernel_args, kernel_args) + - Sets flattened kernel args, kernel arg types and number of those + * - :py:obj:`set_arguments_form_tuple `\ (ty_kernel_args_tuple, ll_kernel_args_tuple) + - Sets flattened kernel args, kernel arg types and number of those + * - :py:obj:`set_dependent_events `\ (dep_events) + - Sets dependent events to the argument list. + * - :py:obj:`set_dependent_events_from_tuple `\ (ty_dependent_events, ll_dependent_events) + - Set's dependent events from tuple represented by LLVM IR. + * - :py:obj:`submit `\ () + - Submits kernel by calling sycl.dpctl_queue_submit_range or + * - :py:obj:`acquire_meminfo_and_submit_release `\ () + - Schedule sycl host task to release nrt meminfo of the arguments used + + + .. rubric:: Members + + .. py:method:: dpexrt() + + Dpex runtime context. + + + .. py:method:: get_queue(exec_queue: dpctl.SyclQueue) -> llvmlite.ir.Instruction + + Allocates memory on the stack to store a DPCTLSyclQueueRef. + + Returns: A LLVM Value storing the pointer to the SYCL queue created + using the filter string for the Python exec_queue (dpctl.SyclQueue). + + + .. py:method:: set_kernel(sycl_kernel_ref: llvmlite.ir.Instruction) + + Sets kernel to the argument list. + + + .. py:method:: set_kernel_from_spirv(kernel_module: SPIRVKernelModule, debug=False) + + Sets kernel to the argument list from the SPIRV bytecode. + + It pastes bytecode as a constant string and create kernel bundle from it + using SYCL API. It caches kernel, so it won't be sent to device second + time. + + + .. py:method:: set_queue(sycl_queue_ref: llvmlite.ir.Instruction) + + Sets queue to the argument list. + + + .. py:method:: set_queue_from_arguments() + + Sets the sycl queue from the first USMNdArray argument provided + earlier. + + + .. py:method:: set_range(global_range: list, local_range: list = None) + + Sets global and local range if provided to the argument list. + + + .. py:method:: set_range_from_indexer(ty_indexer_arg: Union[numba_dpex.core.types.kernel_api.ranges.RangeType, numba_dpex.core.types.kernel_api.ranges.NdRangeType], ll_index_arg: llvmlite.ir.BaseStructType) + + Returns two lists of LLVM IR Values that hold the unboxed extents of + a Python Range or NdRange object. + + + .. py:method:: set_arguments(ty_kernel_args: list[numba.core.types.Type], kernel_args: list[llvmlite.ir.Instruction]) + + Sets flattened kernel args, kernel arg types and number of those + arguments to the argument list. + + + .. py:method:: set_arguments_form_tuple(ty_kernel_args_tuple: numba.core.types.containers.UniTuple, ll_kernel_args_tuple: llvmlite.ir.Instruction) + + Sets flattened kernel args, kernel arg types and number of those + arguments to the argument list based on the arguments stored in tuple. + + + .. py:method:: set_dependent_events(dep_events: list[llvmlite.ir.Instruction]) + + Sets dependent events to the argument list. + + + .. py:method:: set_dependent_events_from_tuple(ty_dependent_events: numba.core.types.containers.UniTuple, ll_dependent_events: llvmlite.ir.Instruction) + + Set's dependent events from tuple represented by LLVM IR. + + :param ll_dependent_events: tuple of numba's data models. + + + .. py:method:: submit() -> llvmlite.ir.Instruction + + Submits kernel by calling sycl.dpctl_queue_submit_range or + sycl.dpctl_queue_submit_ndrange. Must be called after all arguments + set. + + + .. py:method:: acquire_meminfo_and_submit_release() -> llvmlite.ir.Instruction + + Schedule sycl host task to release nrt meminfo of the arguments used + to run job. Use it to keep arguments alive during kernel execution. + + + + +Functions +--------- +.. py:function:: get_queue_from_llvm_values(ctx: numba.core.cpu.CPUContext, builder: llvmlite.ir.builder.IRBuilder, ty_kernel_args: list[numba.core.types.Type], ll_kernel_args: list[llvmlite.ir.Instruction]) + + Get the sycl queue from the first USMNdArray argument. Prior passes + before lowering make sure that compute-follows-data is enforceable + for a specific call to a kernel. As such, at the stage of lowering + the queue from the first USMNdArray argument can be extracted. + + + +Attributes +---------- +.. py:data:: MAX_SIZE_OF_SYCL_RANGE + :value: 3 + + + +.. py:data:: OPEN_CL_OPT_DISABLE_FLAG + :value: '-cl-opt-disable' + + + +.. py:data:: L0_OPT_DISABLE_FLAG + :value: '-g' + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/index.rst.txt new file mode 100644 index 0000000000..7d62050b62 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/index.rst.txt @@ -0,0 +1,61 @@ + +numba_dpex.core.utils.kernel_templates +====================================== + +.. py:module:: numba_dpex.core.utils.kernel_templates + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`RangeKernelTemplate ` + - A template class to generate a numba_dpex.kernel decorated function + + + + +Classes +------- + +.. py:class:: RangeKernelTemplate(kernel_name, kernel_params, kernel_rank, ivar_names, sentinel_name, loop_ranges, param_dict) + + A template class to generate a numba_dpex.kernel decorated function + representing a basic range kernel. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dump_kernel_string `\ () + - Helper to print the kernel function string. + * - :py:obj:`dump_kernel_ir `\ () + - Helper to dump the Numba IR for the RangeKernelTemplate. + + + .. rubric:: Members + + .. py:method:: dump_kernel_string() + + Helper to print the kernel function string. + + + .. py:method:: dump_kernel_ir() + + Helper to dump the Numba IR for the RangeKernelTemplate. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/kernel_template_iface/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/kernel_template_iface/index.rst.txt new file mode 100644 index 0000000000..9fb8fb0827 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/kernel_template_iface/index.rst.txt @@ -0,0 +1,58 @@ + +:orphan: + +numba_dpex.core.utils.kernel_templates.kernel_template_iface +============================================================ + +.. py:module:: numba_dpex.core.utils.kernel_templates.kernel_template_iface + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`KernelTemplateInterface ` + - \- + + + + +Classes +------- + +.. py:class:: KernelTemplateInterface + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dump_kernel_string `\ () + - :summarylabel:`abc` \- + * - :py:obj:`dump_kernel_ir `\ () + - :summarylabel:`abc` \- + + + .. rubric:: Members + + .. py:method:: dump_kernel_string() + :abstractmethod: + + + .. py:method:: dump_kernel_ir() + :abstractmethod: + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/range_kernel_template/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/range_kernel_template/index.rst.txt new file mode 100644 index 0000000000..fe6de021df --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/range_kernel_template/index.rst.txt @@ -0,0 +1,63 @@ + +:orphan: + +numba_dpex.core.utils.kernel_templates.range_kernel_template +============================================================ + +.. py:module:: numba_dpex.core.utils.kernel_templates.range_kernel_template + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`RangeKernelTemplate ` + - A template class to generate a numba_dpex.kernel decorated function + + + + +Classes +------- + +.. py:class:: RangeKernelTemplate(kernel_name, kernel_params, kernel_rank, ivar_names, sentinel_name, loop_ranges, param_dict) + + A template class to generate a numba_dpex.kernel decorated function + representing a basic range kernel. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dump_kernel_string `\ () + - Helper to print the kernel function string. + * - :py:obj:`dump_kernel_ir `\ () + - Helper to dump the Numba IR for the RangeKernelTemplate. + + + .. rubric:: Members + + .. py:method:: dump_kernel_string() + + Helper to print the kernel function string. + + + .. py:method:: dump_kernel_ir() + + Helper to dump the Numba IR for the RangeKernelTemplate. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/reduction_template/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/reduction_template/index.rst.txt new file mode 100644 index 0000000000..613423ea06 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/kernel_templates/reduction_template/index.rst.txt @@ -0,0 +1,105 @@ + +:orphan: + +numba_dpex.core.utils.kernel_templates.reduction_template +========================================================= + +.. py:module:: numba_dpex.core.utils.kernel_templates.reduction_template + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`TreeReduceIntermediateKernelTemplate ` + - The class to build reduction main kernel_txt template and + * - :py:obj:`RemainderReduceIntermediateKernelTemplate ` + - The class to build reduction remainder kernel_txt template and + + + + +Classes +------- + +.. py:class:: TreeReduceIntermediateKernelTemplate(kernel_name, kernel_params, ivar_names, sentinel_name, loop_ranges, param_dict, parfor_dim, redvars, parfor_args, parfor_reddict, redvars_dict, typemap, work_group_size) + + Bases: :py:obj:`numba_dpex.core.utils.kernel_templates.kernel_template_iface.KernelTemplateInterface` + + The class to build reduction main kernel_txt template and + compiled Numba functionIR. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dump_kernel_string `\ () + - Helper to print the kernel function string. + * - :py:obj:`dump_kernel_ir `\ () + - Helper to dump the Numba IR for a + + + .. rubric:: Members + + .. py:method:: dump_kernel_string() + + Helper to print the kernel function string. + + + .. py:method:: dump_kernel_ir() + + Helper to dump the Numba IR for a + TreeReduceIntermediateKernelTemplate. + + + + +.. py:class:: RemainderReduceIntermediateKernelTemplate(kernel_name, kernel_params, sentinel_name, redvars, parfor_reddict, redvars_dict, typemap, legal_loop_indices, global_size_var_name, global_size_mod_var_name, partial_sum_size_var_name, partial_sum_var_name, final_sum_var_name, reductionKernelVar) + + Bases: :py:obj:`numba_dpex.core.utils.kernel_templates.kernel_template_iface.KernelTemplateInterface` + + The class to build reduction remainder kernel_txt template and + compiled Numba functionIR. + + + .. rubric:: Overview + + + .. list-table:: Methods + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`dump_kernel_string `\ () + - Helper to print the kernel function string. + * - :py:obj:`dump_kernel_ir `\ () + - Helper to dump the Numba IR for the + + + .. rubric:: Members + + .. py:method:: dump_kernel_string() + + Helper to print the kernel function string. + + + .. py:method:: dump_kernel_ir() + + Helper to dump the Numba IR for the + RemainderReduceIntermediateKernelTemplate. + + + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/core/utils/suai_helper/index.rst.txt b/dev/_sources/autoapi/numba_dpex/core/utils/suai_helper/index.rst.txt new file mode 100644 index 0000000000..4d4afe4729 --- /dev/null +++ b/dev/_sources/autoapi/numba_dpex/core/utils/suai_helper/index.rst.txt @@ -0,0 +1,63 @@ + +:orphan: + +numba_dpex.core.utils.suai_helper +================================= + +.. py:module:: numba_dpex.core.utils.suai_helper + + +Overview +-------- + +.. list-table:: Classes + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`SyclUSMArrayInterface ` + - Stores as attributes the information extracted from a + + +.. list-table:: Function + :header-rows: 0 + :widths: auto + :class: summarytable + + * - :py:obj:`get_info_from_suai `\ (obj) + - Extracts the metadata of an object of type UsmNdArray using the objects + + + +Classes +------- + +.. py:class:: SyclUSMArrayInterface(data, writable, size, shape, dimensions, itemsize, strides, dtype, usm_type, device, queue) + + Stores as attributes the information extracted from a + __sycl_usm_array_interface__ dictionary as defined by dpctl.memory.Memory* + classes. + + + + +Functions +--------- +.. py:function:: get_info_from_suai(obj) + + Extracts the metadata of an object of type UsmNdArray using the objects + __sycl_usm_array_interface__ (SUAI) attribute. + + The ``dpctl.memory.as_usm_memory`` function converts the array-like + object into a dpctl.memory.USMMemory object. Using the ``as_usm_memory`` + is an implicit way to verify if the array-like object is a legal + SYCL USM memory back Python object that can be passed to a dpex kernel. + + :param obj: array-like object with a SUAI attribute. + + :returns: A SyclUSMArrayInterface object + + + + + diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_fence_overloads/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_fence_overloads/index.rst.txt deleted file mode 100644 index 7c00575a43..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_fence_overloads/index.rst.txt +++ /dev/null @@ -1,44 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads._atomic_fence_overloads -============================================================================= - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads._atomic_fence_overloads - -.. autoapi-nested-parse:: - - Provides overloads for functions included in kernel_api.atomic_fence - that generate dpcpp SPIR-V LLVM IR intrinsic function calls. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`ol_atomic_fence `\ (memory_order, memory_scope) - - SPIR-V overload for - - - - -Functions ---------- -.. py:function:: ol_atomic_fence(memory_order, memory_scope) - - SPIR-V overload for - :meth:`numba_dpex.kernel_api.atomic_fence`. - - Generates the same LLVM IR instruction as DPC++ for the SYCL - `atomic_fence` function. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads/index.rst.txt deleted file mode 100644 index e8159a1230..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads/index.rst.txt +++ /dev/null @@ -1,193 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads._atomic_ref_overloads -=========================================================================== - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads._atomic_ref_overloads - -.. autoapi-nested-parse:: - - Implements the SPIR-V overloads for the kernel_api.AtomicRef class methods. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`ol_atomic_ref `\ (ref, index, memory_order, memory_scope, address_space) - - Overload of the constructor for the class - * - :py:obj:`ol_fetch_add `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_add`. - * - :py:obj:`ol_fetch_sub `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_sub`. - * - :py:obj:`ol_fetch_min `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_min`. - * - :py:obj:`ol_fetch_max `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_max`. - * - :py:obj:`ol_fetch_and `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_and`. - * - :py:obj:`ol_fetch_or `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_or`. - * - :py:obj:`ol_fetch_xor `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_xor`. - * - :py:obj:`ol_load `\ (atomic_ref) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.load`. - * - :py:obj:`ol_store `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.store`. - * - :py:obj:`ol_exchange `\ (atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.exchange`. - * - :py:obj:`ol_compare_exchange `\ (atomic_ref, expected_ref, desired, expected_idx) - - SPIR-V overload for - - - - -Functions ---------- -.. py:function:: ol_atomic_ref(ref, index, memory_order=MemoryOrder.RELAXED, memory_scope=MemoryScope.DEVICE, address_space=AddressSpace.GLOBAL) - - Overload of the constructor for the class - class:`numba_dpex.kernel_api.AtomicRef`. - - :raises errors.TypingError: If the `ref` argument is not a UsmNdArray type. - :raises errors.TypingError: If the dtype of the `ref` is not supported in an - :raises AtomicRef.: - :raises errors.TypingError: If the device does not support atomic operations on - :raises the dtype of the `ref`.: - :raises errors.TypingError: If the `memory_order`, `address_type`, or - :raises memory_scope: - :raises errors.TypingError: If the `address_space` argument is different from - :raises the address space attribute of the `ref` argument.: - :raises errors.TypingError: If the address space is PRIVATE. - - -.. py:function:: ol_fetch_add(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_add`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::fetch_add` function. - - :raises TypingError: When the dtype of the aggregator value does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_fetch_sub(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_sub`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::fetch_sub` function. - - :raises TypingError: When the dtype of the aggregator value does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_fetch_min(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_min`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::fetch_min` function. - - :raises TypingError: When the dtype of the aggregator value does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_fetch_max(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_max`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::fetch_max` function. - - :raises TypingError: When the dtype of the aggregator value does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_fetch_and(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_and`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::fetch_and` function. - - :raises TypingError: When the dtype of the aggregator value does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_fetch_or(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_or`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::fetch_or` function. - - :raises TypingError: When the dtype of the aggregator value does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_fetch_xor(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_xor`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::fetch_xor` function. - - :raises TypingError: When the dtype of the aggregator value does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_load(atomic_ref) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.load`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::load` function. - - - -.. py:function:: ol_store(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.store`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::store` function. - - :raises TypingError: When the dtype of the value stored does not match the - :raises dtype of the AtomicRef type.: - - -.. py:function:: ol_exchange(atomic_ref, val) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.exchange`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::exchange` function. - - :raises TypingError: When the dtype of the value passed to `exchange` - :raises does not match the dtype of the AtomicRef type.: - - -.. py:function:: ol_compare_exchange(atomic_ref, expected_ref, desired, expected_idx=0) - - SPIR-V overload for - :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.compare_exchange`. - - Generates the same LLVM IR instruction as dpcpp for the - `atomic_ref::compare_exchange_strong` function. - - :raises TypingError: When the dtype of the value passed to `compare_exchange` - :raises does not match the dtype of the AtomicRef type.: - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads/index.rst.txt deleted file mode 100644 index 4e15e1695c..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads/index.rst.txt +++ /dev/null @@ -1,55 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads._group_barrier_overloads -============================================================================== - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads._group_barrier_overloads - -.. autoapi-nested-parse:: - - Provides overloads for functions included in kernel_api.barrier that - generate dpcpp SPIR-V LLVM IR intrinsic function calls. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`ol_group_barrier `\ (group, fence_scope) - - SPIR-V overload for - - - - -Functions ---------- -.. py:function:: ol_group_barrier(group, fence_scope=MemoryScope.WORK_GROUP) - - SPIR-V overload for - :meth:`numba_dpex.kernel_api.group_barrier`. - - Generates the same LLVM IR instruction as DPC++ for the SYCL - `group_barrier` function. - - Per SYCL spec, group_barrier must perform both control barrier and memory - fence operations. Hence, group_barrier requires two scopes and one memory - consistency specification as its three arguments. - - mem_scope - scope of any memory consistency operations that are performed by - the barrier. By default, mem_scope is set to `work_group`. - exec_scope - scope that determines the set of work-items that synchronize at - barrier. Set to `work_group` for group_barrier always. - spirv_memory_semantics_mask - Based on SYCL implementation. Always set to - use sequential consistency memory order. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads/index.rst.txt deleted file mode 100644 index 3d53c1a43e..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads/index.rst.txt +++ /dev/null @@ -1,94 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads._index_space_id_overloads -=============================================================================== - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads._index_space_id_overloads - -.. autoapi-nested-parse:: - - Implements the SPIR-V overloads for the kernel_api.items class methods. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`spirv_name `\ (name) - - Converts name to spirv name by adding __spirv_ prefix. - * - :py:obj:`declare_spirv_const `\ (builder, name) - - Declares global external spirv constant - * - :py:obj:`generate_index_overload `\ (_type, _intrinsic) - - Generates overload for the index method that generates specific IR from - * - :py:obj:`register_index_const_methods `\ () - - Register indexing related methods that can be defined as spirv const. - * - :py:obj:`ol_nd_item_get_group `\ (nd_item) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.NdItem.get_group`. - * - :py:obj:`ol_nd_item_dimensions `\ (item) - - SPIR-V overload for :meth:`numba_dpex.kernel_api..dimensions`. - * - :py:obj:`register_jitable_method `\ (type_, method) - - Register a regular python method that can be executed by the - - - - -Functions ---------- -.. py:function:: spirv_name(name: str) - - Converts name to spirv name by adding __spirv_ prefix. - - -.. py:function:: declare_spirv_const(builder: llvmlite.ir.IRBuilder, name: str) - - Declares global external spirv constant - - -.. py:function:: generate_index_overload(_type, _intrinsic) - - Generates overload for the index method that generates specific IR from - provided intrinsic. - - -.. py:function:: register_index_const_methods() - - Register indexing related methods that can be defined as spirv const. - - -.. py:function:: ol_nd_item_get_group(nd_item) - - SPIR-V overload for :meth:`numba_dpex.kernel_api.NdItem.get_group`. - - Generates the same LLVM IR instruction as dpcpp for the - `sycl::nd_item::get_group` function. - - :raises TypingError: When argument is not NdItem. - - -.. py:function:: ol_nd_item_dimensions(item) - - SPIR-V overload for :meth:`numba_dpex.kernel_api..dimensions`. - - Generates the same LLVM IR instruction as dpcpp for the - `sycl::::dimensions` attribute. - - -.. py:function:: register_jitable_method(type_, method) - - Register a regular python method that can be executed by the - python interpreter and can be compiled into a nopython - function when referenced by other jit'ed functions. - - Same as register_jitable, but for methods with no arguments. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_private_array_overloads/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_private_array_overloads/index.rst.txt deleted file mode 100644 index a3039c2cc6..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_private_array_overloads/index.rst.txt +++ /dev/null @@ -1,52 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads._private_array_overloads -============================================================================== - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads._private_array_overloads - -.. autoapi-nested-parse:: - - Implements the SPIR-V overloads for the kernel_api.PrivateArray class. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`type_interval `\ (context) - - Sets type of the constructor for the class - * - :py:obj:`dpex_private_array_lower `\ (context, builder, sig, args) - - Implements lower for the class:`numba_dpex.kernel_api.PrivateArray` - - - - -Functions ---------- -.. py:function:: type_interval(context) - - Sets type of the constructor for the class - class:`numba_dpex.kernel_api.PrivateArray`. - - :raises errors.TypingError: If the shape argument is not a shape compatible - type. - :raises errors.TypingError: If the dtype argument is not a dtype compatible - type. - - -.. py:function:: dpex_private_array_lower(context: numba_dpex.experimental.target.DpexExpKernelTypingContext, builder: llvmlite.ir.builder.IRBuilder, sig: numba.core.typing.templates.Signature, args: list[llvmlite.ir.Value]) - - Implements lower for the class:`numba_dpex.kernel_api.PrivateArray` - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_registry/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_registry/index.rst.txt deleted file mode 100644 index f7061cc4a6..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_registry/index.rst.txt +++ /dev/null @@ -1,32 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads._registry -=============================================================== - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads._registry - -.. autoapi-nested-parse:: - - Implements the SPIR-V overloads for the kernel_api.PrivateArray class. - - - -Overview --------- - - - - -Attributes ----------- -.. py:data:: registry - - - -.. py:data:: lower - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper/index.rst.txt deleted file mode 100644 index 391d66826a..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper/index.rst.txt +++ /dev/null @@ -1,74 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads._spv_atomic_inst_helper -============================================================================= - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads._spv_atomic_inst_helper - -.. autoapi-nested-parse:: - - Helper module to generate LLVM IR SPIR-V intrinsic calls. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`get_atomic_inst_name `\ (atomic_inst, atomic_ref_dtype) - - Returns a string corresponding to the LLVM IR intrinsic function - * - :py:obj:`get_memory_semantics_mask `\ (memory_order) - - Equivalent of the getMemorySemanticsMask function in dpcpp's - * - :py:obj:`get_scope `\ (memory_scope) - - Translates a memory scope enum value to an equivalent SPIR-V memory scope - - - - -Functions ---------- -.. py:function:: get_atomic_inst_name(atomic_inst: str, atomic_ref_dtype: numba.core.types.Type) -> str - - Returns a string corresponding to the LLVM IR intrinsic function - generated for a particular atomic operation for a specific data type. - - :param atomic_inst: The string name of an atomic operation to look up. - :type atomic_inst: str - :param atomic_ref_dtype: The Numba type object - :type atomic_ref_dtype: numba.core.types.Type - :param corresponding to the dtype of the object stored in an AtomicRef.: - - :raises ValueError: If the atomic operation is not found in the - :raises `_spv_atomic_instructions_map`` dictionary: - :raises ValueError: If no instruction is found for a combination of the - :raises atomic operation and the dtype in the ``_spv_atomic_instructions_map``: - :raises dictionary.: - - :returns: A string corresponding to the LLVM IR intrinsic that should be - generated for the atomic operation for the specified dtype. - :rtype: str - - -.. py:function:: get_memory_semantics_mask(memory_order) - - Equivalent of the getMemorySemanticsMask function in dpcpp's - sycl/include/sycl/detail/spirv.hpp. Translates SYCL memory order to a - SPIR-V memory semantics mask. - - - -.. py:function:: get_scope(memory_scope) - - Translates a memory scope enum value to an equivalent SPIR-V memory scope - value. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/index.rst.txt deleted file mode 100644 index c09a73fe9f..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/index.rst.txt +++ /dev/null @@ -1,9 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads -===================================================== - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/spv_fn_declarations/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/spv_fn_declarations/index.rst.txt deleted file mode 100644 index 895fbc3a0c..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/spv_fn_declarations/index.rst.txt +++ /dev/null @@ -1,81 +0,0 @@ - -:orphan: - -numba_dpex.experimental._kernel_dpcpp_spirv_overloads.spv_fn_declarations -========================================================================= - -.. py:module:: numba_dpex.experimental._kernel_dpcpp_spirv_overloads.spv_fn_declarations - -.. autoapi-nested-parse:: - - Implements a set of helper functions to generate the LLVM IR for SPIR-V - functions and their use inside an LLVM module. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`get_or_insert_atomic_load_fn `\ (context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicLoad call into the - * - :py:obj:`get_or_insert_spv_atomic_store_fn `\ (context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicStore call into the - * - :py:obj:`get_or_insert_spv_atomic_exchange_fn `\ (context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicExchange call into the - * - :py:obj:`get_or_insert_spv_atomic_compare_exchange_fn `\ (context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicCompareExchange call into the - * - :py:obj:`get_or_insert_spv_group_barrier_fn `\ (module) - - Gets or inserts a declaration for a __spirv_ControlBarrier call into the - * - :py:obj:`get_or_insert_spv_atomic_fence_fn `\ (module) - - Gets or inserts a declaration for a __spirv_MemoryBarrier call into the - - - - -Functions ---------- -.. py:function:: get_or_insert_atomic_load_fn(context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicLoad call into the - specified LLVM IR module. - - -.. py:function:: get_or_insert_spv_atomic_store_fn(context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicStore call into the - specified LLVM IR module. - - -.. py:function:: get_or_insert_spv_atomic_exchange_fn(context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicExchange call into the - specified LLVM IR module. - - -.. py:function:: get_or_insert_spv_atomic_compare_exchange_fn(context, module, atomic_ref_ty) - - Gets or inserts a declaration for a __spirv_AtomicCompareExchange call into the - specified LLVM IR module. - - -.. py:function:: get_or_insert_spv_group_barrier_fn(module) - - Gets or inserts a declaration for a __spirv_ControlBarrier call into the - specified LLVM IR module. - - -.. py:function:: get_or_insert_spv_atomic_fence_fn(module) - - Gets or inserts a declaration for a __spirv_MemoryBarrier call into the - specified LLVM IR module. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/decorators/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/decorators/index.rst.txt deleted file mode 100644 index 6d1f8cf37a..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/decorators/index.rst.txt +++ /dev/null @@ -1,69 +0,0 @@ - -:orphan: - -numba_dpex.experimental.decorators -================================== - -.. py:module:: numba_dpex.experimental.decorators - -.. autoapi-nested-parse:: - - The set of experimental decorators provided by numba_dpex that are not yet - ready to move to numba_dpex.core. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`kernel `\ (func_or_sig, \*\*options) - - A decorator to define a kernel function. - * - :py:obj:`device_func `\ (func_or_sig, \*\*options) - - Generates a function with a device-only calling convention, e.g., - - - - -Functions ---------- -.. py:function:: kernel(func_or_sig=None, **options) - - A decorator to define a kernel function. - - A kernel function is conceptually equivalent to a SYCL kernel function, and - gets compiled into either an OpenCL or a LevelZero SPIR-V binary kernel. - A kernel decorated Python function has the following restrictions: - - * The function can not return any value. - * All array arguments passed to a kernel should adhere to compute - follows data programming model. - - -.. py:function:: device_func(func_or_sig=None, **options) - - Generates a function with a device-only calling convention, e.g., - spir_func for SPIR-V based devices. - - The decorator is used to compile overloads in the DpexKernelTarget and - users should use the decorator to define functions that are only callable - from inside another device_func or a kernel. - - A device_func is not compiled down to device binary IR and instead left as - LLVM IR. It is done so that the function can be inlined fully into the - kernel module from where it is used at the LLVM level, leading to more - optimization opportunities. - - :returns: A KernelDispatcher instance with the - _compilation_mode option set to DEVICE_FUNC. - :rtype: KernelDispatcher - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/index.rst.txt deleted file mode 100644 index 93a753aabe..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/index.rst.txt +++ /dev/null @@ -1,119 +0,0 @@ - -numba_dpex.experimental -======================= - -.. py:module:: numba_dpex.experimental - -.. autoapi-nested-parse:: - - Contains experimental features that are meant as engineering preview and not - yet production ready. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`device_func `\ (func_or_sig, \*\*options) - - Generates a function with a device-only calling convention, e.g., - * - :py:obj:`kernel `\ (func_or_sig, \*\*options) - - A decorator to define a kernel function. - * - :py:obj:`call_kernel `\ (kernel_fn, index_space, \*kernel_args) - - Compiles and synchronously executes a kernel function. - * - :py:obj:`call_kernel_async `\ (kernel_fn, index_space, dependent_events, \*kernel_args) - - Compiles and asynchronously executes a kernel function. - - - - -Functions ---------- -.. py:function:: device_func(func_or_sig=None, **options) - - Generates a function with a device-only calling convention, e.g., - spir_func for SPIR-V based devices. - - The decorator is used to compile overloads in the DpexKernelTarget and - users should use the decorator to define functions that are only callable - from inside another device_func or a kernel. - - A device_func is not compiled down to device binary IR and instead left as - LLVM IR. It is done so that the function can be inlined fully into the - kernel module from where it is used at the LLVM level, leading to more - optimization opportunities. - - :returns: A KernelDispatcher instance with the - _compilation_mode option set to DEVICE_FUNC. - :rtype: KernelDispatcher - - -.. py:function:: kernel(func_or_sig=None, **options) - - A decorator to define a kernel function. - - A kernel function is conceptually equivalent to a SYCL kernel function, and - gets compiled into either an OpenCL or a LevelZero SPIR-V binary kernel. - A kernel decorated Python function has the following restrictions: - - * The function can not return any value. - * All array arguments passed to a kernel should adhere to compute - follows data programming model. - - -.. py:function:: call_kernel(kernel_fn, index_space, *kernel_args) -> None - - Compiles and synchronously executes a kernel function. - - Kernel execution happens in synchronous way, so the main thread will be - blocked till the kernel done execution. - - :param kernel_fn: A - :func:`numba_dpex.experimental.kernel` decorated function that is - compiled to a ``KernelDispatcher``. - :type kernel_fn: numba_dpex.experimental.KernelDispatcher - :param index_space: A Range or NdRange type object that - specifies the index space for the kernel. - :type index_space: Range | NdRange - :param kernel_args: List of objects that are passed to the numba_dpex.kernel - decorated function. - - -.. py:function:: call_kernel_async(kernel_fn, index_space, dependent_events: list[dpctl.SyclEvent], *kernel_args) -> tuple[dpctl.SyclEvent, dpctl.SyclEvent] - - Compiles and asynchronously executes a kernel function. - - Calls a :func:`numba_dpex.experimental.kernel` decorated function - asynchronously from CPython or from a :func:`numba_dpex.dpjit` function. As - the kernel execution happens asynchronously, so the main thread will not be - blocked till the kernel done execution. Instead the function returns back to - caller a handle for an *event* to track kernel execution. It is a user's - responsibility to properly track kernel execution completion and not use any - data that may still be used by the kernel prior to the kernel's completion. - - :param kernel_fn: A - :func:`numba_dpex.experimental.kernel` decorated function that is - compiled to a ``KernelDispatcher``. - :type kernel_fn: KernelDispatcher - :param index_space: A Range or NdRange type object that - specifies the index space for the kernel. - :type index_space: Range | NdRange - :param kernel_args: List of objects that are passed to the numba_dpex.kernel - decorated function. - - :returns: A pair of ``dpctl.SyclEvent`` objects. The pair of events constitute of - a host task and an event associated with the kernel execution. The event - associated with the kernel execution indicates the execution status of - the submitted kernel function. The host task manages the lifetime of any - PyObject passed in as a kernel argument and automatically decrements the - reference count of the object on kernel execution completion. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/models/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/models/index.rst.txt deleted file mode 100644 index 261f1ca0c4..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/models/index.rst.txt +++ /dev/null @@ -1,97 +0,0 @@ - -:orphan: - -numba_dpex.experimental.models -============================== - -.. py:module:: numba_dpex.experimental.models - -.. autoapi-nested-parse:: - - Provides the Numba data models for the numba_dpex types introduced in the - numba_dpex.experimental module. - - - -Overview --------- - -.. list-table:: Classes - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`AtomicRefModel ` - - Data model for AtomicRefType. - * - :py:obj:`EmptyStructModel ` - - Data model that does not take space. Intended to be used with types that - * - :py:obj:`DpctlMDLocalAccessorModel ` - - Data model to represent DpctlMDLocalAccessorType. - * - :py:obj:`LocalAccessorModel ` - - Data model for the LocalAccessor type when used in a host-only function. - - - -.. list-table:: Attributes - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`exp_dmm ` - - \- - - -Classes -------- - -.. py:class:: AtomicRefModel(dmm, fe_type) - - Bases: :py:obj:`numba.core.datamodel.models.StructModel` - - Data model for AtomicRefType. - - - - -.. py:class:: EmptyStructModel(dmm, fe_type) - - Bases: :py:obj:`numba.core.datamodel.models.StructModel` - - Data model that does not take space. Intended to be used with types that - are presented only at typing stage and not represented physically. - - - - -.. py:class:: DpctlMDLocalAccessorModel(dmm, fe_type) - - Bases: :py:obj:`numba.core.datamodel.models.StructModel` - - Data model to represent DpctlMDLocalAccessorType. - - Must be the same structure as - dpctl/syclinterface/dpctl_sycl_queue_interface.h::MDLocalAccessor. - - Structure intended to be used only on host side of the kernel call. - - - - -.. py:class:: LocalAccessorModel(dmm, fe_type) - - Bases: :py:obj:`numba.core.datamodel.models.StructModel` - - Data model for the LocalAccessor type when used in a host-only function. - - - - - -Attributes ----------- -.. py:data:: exp_dmm - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/target/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/target/index.rst.txt deleted file mode 100644 index cb0749ec73..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/target/index.rst.txt +++ /dev/null @@ -1,146 +0,0 @@ - -:orphan: - -numba_dpex.experimental.target -============================== - -.. py:module:: numba_dpex.experimental.target - -.. autoapi-nested-parse:: - - A new target descriptor that includes experimental features that should - eventually move into the numba_dpex.core. - - - -Overview --------- - -.. list-table:: Classes - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`SyclDeviceExp ` - - Mark the hardware target as SYCL Device. - * - :py:obj:`DpexExpKernelTypingContext ` - - Experimental typing context class extending the DpexKernelTypingContext - * - :py:obj:`DpexExpKernelTargetContext ` - - Experimental target context class extending the DpexKernelTargetContext - * - :py:obj:`DpexExpKernelTarget ` - - Implements a target descriptor for numba_dpex.kernel decorated functions. - - - -.. list-table:: Attributes - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`DPEX_KERNEL_EXP_TARGET_NAME ` - - \- - * - :py:obj:`dpex_exp_kernel_target ` - - \- - - -Classes -------- - -.. py:class:: SyclDeviceExp - - Bases: :py:obj:`numba.core.target_extension.GPU` - - Mark the hardware target as SYCL Device. - - - - -.. py:class:: DpexExpKernelTypingContext - - Bases: :py:obj:`numba_dpex.kernel_api_impl.spirv.target.SPIRVTypingContext` - - Experimental typing context class extending the DpexKernelTypingContext - by overriding super class functions for new experimental types. - - A new experimental type may require updating type inference for that type - when it is used as an argument, value or attribute in a JIT compiled - function. All such experimental functionality should be added here till they - are stable enough to be migrated to DpexKernelTypingContext. - - - - -.. py:class:: DpexExpKernelTargetContext(typingctx, target=DPEX_KERNEL_EXP_TARGET_NAME) - - Bases: :py:obj:`numba_dpex.kernel_api_impl.spirv.target.SPIRVTargetContext` - - Experimental target context class extending the DpexKernelTargetContext - by overriding super class functions for new experimental types. - - A new experimental type may require specific ways for handling the lowering - to LLVM IR. All such experimental functionality should be added here till - they are stable enough to be migrated to DpexKernelTargetContext. - - - .. rubric:: Overview - - .. list-table:: Attributes - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`allow_dynamic_globals ` - - \- - - - - .. rubric:: Members - - .. py:attribute:: allow_dynamic_globals - :value: True - - - - - -.. py:class:: DpexExpKernelTarget(target_name) - - Bases: :py:obj:`numba.core.descriptors.TargetDescriptor` - - Implements a target descriptor for numba_dpex.kernel decorated functions. - - - .. rubric:: Overview - - .. list-table:: Attributes - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`options ` - - \- - - - - .. rubric:: Members - - .. py:attribute:: options - - - - - - -Attributes ----------- -.. py:data:: DPEX_KERNEL_EXP_TARGET_NAME - :value: 'dpex_kernel_exp' - - - -.. py:data:: dpex_exp_kernel_target - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/testing/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/testing/index.rst.txt deleted file mode 100644 index 44ae829669..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/testing/index.rst.txt +++ /dev/null @@ -1,39 +0,0 @@ - -:orphan: - -numba_dpex.experimental.testing -=============================== - -.. py:module:: numba_dpex.experimental.testing - -.. autoapi-nested-parse:: - - Tools for testing, not intended for regular use. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`kernel_cache_size `\ () - - Returns kernel cache size. - - - - -Functions ---------- -.. py:function:: kernel_cache_size() -> int - - Returns kernel cache size. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/typeof/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/typeof/index.rst.txt deleted file mode 100644 index 64a9e7b44b..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/typeof/index.rst.txt +++ /dev/null @@ -1,103 +0,0 @@ - -:orphan: - -numba_dpex.experimental.typeof -============================== - -.. py:module:: numba_dpex.experimental.typeof - -.. autoapi-nested-parse:: - - Adds typeof implementations to Numba registry for all numba-dpex experimental - types. - - - -Overview --------- - - -.. list-table:: Function - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`typeof_atomic_ref `\ (val, ctx) - - Returns a ``numba_dpex.experimental.dpctpp_types.AtomicRefType`` - * - :py:obj:`typeof_group `\ (val, c) - - Registers the type inference implementation function for a - * - :py:obj:`typeof_item `\ (val, c) - - Registers the type inference implementation function for a - * - :py:obj:`typeof_nditem `\ (val, c) - - Registers the type inference implementation function for a - * - :py:obj:`typeof_local_accessor `\ (val, c) - - Returns a ``numba_dpex.experimental.dpctpp_types.LocalAccessorType`` - - - - -Functions ---------- -.. py:function:: typeof_atomic_ref(val: numba_dpex.kernel_api.AtomicRef, ctx) -> numba_dpex.core.types.kernel_api.atomic_ref.AtomicRefType - - Returns a ``numba_dpex.experimental.dpctpp_types.AtomicRefType`` - instance for a Python AtomicRef object. - - :param val: Instance of the AtomicRef type. - :type val: AtomicRef - :param ctx: Numba typing context used for type inference. - - Returns: AtomicRefType object corresponding to the AtomicRef object. - - - -.. py:function:: typeof_group(val: numba_dpex.kernel_api.Group, c) - - Registers the type inference implementation function for a - numba_dpex.kernel_api.Group PyObject. - - :param val: An instance of numba_dpex.kernel_api.Group. - :param c: Unused argument used to be consistent with Numba API. - - Returns: A numba_dpex.experimental.core.types.kernel_api.items.GroupType - instance. - - -.. py:function:: typeof_item(val: numba_dpex.kernel_api.Item, c) - - Registers the type inference implementation function for a - numba_dpex.kernel_api.Item PyObject. - - :param val: An instance of numba_dpex.kernel_api.Item. - :param c: Unused argument used to be consistent with Numba API. - - Returns: A numba_dpex.experimental.core.types.kernel_api.items.ItemType - instance. - - -.. py:function:: typeof_nditem(val: numba_dpex.kernel_api.NdItem, c) - - Registers the type inference implementation function for a - numba_dpex.kernel_api.NdItem PyObject. - - :param val: An instance of numba_dpex.kernel_api.NdItem. - :param c: Unused argument used to be consistent with Numba API. - - Returns: A numba_dpex.experimental.core.types.kernel_api.items.NdItemType - instance. - - -.. py:function:: typeof_local_accessor(val: numba_dpex.kernel_api.LocalAccessor, c) -> numba_dpex.core.types.kernel_api.local_accessor.LocalAccessorType - - Returns a ``numba_dpex.experimental.dpctpp_types.LocalAccessorType`` - instance for a Python LocalAccessor object. - :param val: Instance of the LocalAccessor type. - :type val: LocalAccessor - :param c: Numba typing context used for type inference. - - Returns: LocalAccessorType object corresponding to the LocalAccessor object. - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/experimental/types/index.rst.txt b/dev/_sources/autoapi/numba_dpex/experimental/types/index.rst.txt deleted file mode 100644 index 109e6d113d..0000000000 --- a/dev/_sources/autoapi/numba_dpex/experimental/types/index.rst.txt +++ /dev/null @@ -1,61 +0,0 @@ - -:orphan: - -numba_dpex.experimental.types -============================= - -.. py:module:: numba_dpex.experimental.types - -.. autoapi-nested-parse:: - - Experimental types that will eventually move to numba_dpex.core.types - - - -Overview --------- - -.. list-table:: Classes - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`KernelDispatcherType ` - - The type of KernelDispatcher dispatchers - - - - -Classes -------- - -.. py:class:: KernelDispatcherType(dispatcher) - - Bases: :py:obj:`numba.core.types.Dispatcher` - - The type of KernelDispatcher dispatchers - - - .. rubric:: Overview - - - .. list-table:: Methods - :header-rows: 0 - :widths: auto - :class: summarytable - - * - :py:obj:`cast_python_value `\ (args) - - :summarylabel:`abc` \- - - - .. rubric:: Members - - .. py:method:: cast_python_value(args) - :abstractmethod: - - - - - - - diff --git a/dev/_sources/autoapi/numba_dpex/kernel_api/atomic_ref/index.rst.txt b/dev/_sources/autoapi/numba_dpex/kernel_api/atomic_ref/index.rst.txt index bce32ee955..e8fd30d379 100644 --- a/dev/_sources/autoapi/numba_dpex/kernel_api/atomic_ref/index.rst.txt +++ b/dev/_sources/autoapi/numba_dpex/kernel_api/atomic_ref/index.rst.txt @@ -30,7 +30,7 @@ Overview Classes ------- -.. py:class:: AtomicRef(ref, index, memory_order=MemoryOrder.RELAXED, memory_scope=MemoryScope.DEVICE, address_space=AddressSpace.GLOBAL) +.. py:class:: AtomicRef(ref, index, memory_order=MemoryOrder.RELAXED, memory_scope=MemoryScope.DEVICE, address_space=None) Analogue to the ``sycl::atomic_ref`` class. diff --git a/dev/_sources/autoapi/numba_dpex/kernel_api/index.rst.txt b/dev/_sources/autoapi/numba_dpex/kernel_api/index.rst.txt index 7ab0c871cd..cff30a5e0c 100644 --- a/dev/_sources/autoapi/numba_dpex/kernel_api/index.rst.txt +++ b/dev/_sources/autoapi/numba_dpex/kernel_api/index.rst.txt @@ -62,7 +62,7 @@ Overview Classes ------- -.. py:class:: AtomicRef(ref, index, memory_order=MemoryOrder.RELAXED, memory_scope=MemoryScope.DEVICE, address_space=AddressSpace.GLOBAL) +.. py:class:: AtomicRef(ref, index, memory_order=MemoryOrder.RELAXED, memory_scope=MemoryScope.DEVICE, address_space=None) Analogue to the ``sycl::atomic_ref`` class. diff --git a/dev/autoapi/index.html b/dev/autoapi/index.html index 64766b4dd2..166d64dca6 100644 --- a/dev/autoapi/index.html +++ b/dev/autoapi/index.html @@ -195,8 +195,8 @@
  • API Reference
  • Experimental Features
    • @@ -251,8 +251,8 @@

    API Reference