Move and deprecate dpctl program

.flake8

@@ -22,7 +22,7 @@ per-file-ignores =
     dpctl/_sycl_queue.pyx: E999, E225, E226, E227
     dpctl/_sycl_queue_manager.pyx: E999, E225
     dpctl/memory/_memory.pyx: E999, E225, E226, E227
-    dpctl/program/_program.pyx: E999, E225, E226, E227
+    dpctl/compiler/_compiler.pyx: E999, E225, E226, E227
     dpctl/tests/_cython_api.pyx: E999, E225, E227, E402
     dpctl/utils/_onetrace_context.py: E501, W505
     examples/cython/sycl_buffer/syclbuffer/_syclbuffer.pyx: E999, E225, E402

docs/doc_sources/api_reference/dpctl/compiler.rst

@@ -0,0 +1,38 @@
+.. _dpctl_compiler_pyapi:
+
+:py:mod:`dpctl.compiler`
+========================
+
+:py:mod:`dpctl.compiler` provides a way to create a SYCL kernel
+from either an OpenCL* program source code represented as a string
+or a SPIR-V binary file.
+
+It implements creation of interoperability
+``sycl::kernel_bundle<sycl::bundle_state_executable>`` (a collection of kernels),
+as well as creation of individual ``sycl::kernel``, suitable for submission for
+execution via :py:meth:`dpctl.SyclQueue.submit`.
+
+.. py:module:: dpctl.compiler
+
+.. currentmodule:: dpctl.compiler
+
+.. autosummary::
+    :toctree: generated
+    :nosignatures:
+
+    create_kernel_bundle_from_source
+    create_kernel_bundle_from_spirv
+
+.. autosummary::
+    :toctree: generated
+    :nosignatures:
+
+    SyclKernelBundle
+    SyclKernel
+    SpecializationConstant
+
+.. autosummary::
+    :toctree: generated
+    :nosignatures:
+
+    SyclKernelBundleCompilationError

docs/doc_sources/api_reference/dpctl/index.rst

@@ -14,8 +14,10 @@
 
     * - :py:mod:`dpctl.memory`
       - Unified Shared Memory operations
-    * - :py:mod:`dpctl.program`
+    * - :py:mod:`dpctl.compiler`
       - Support for working with SYCL kernels
+    * - :py:mod:`dpctl.program`
+      - (deprecated, use :py:mod:`dpctl.compiler`)
     * - :py:mod:`dpctl.utils`
       - A collection of utility functions
 

docs/doc_sources/api_reference/dpctl/program.rst

@@ -1,7 +1,10 @@
 .. _dpctl_program_pyapi:
 
-:py:mod:`dpctl.program`
-=======================
+:py:mod:`dpctl.program` (deprecated)
+=====================================
+
+.. deprecated::
+   :py:mod:`dpctl.program` is deprecated. Use :py:mod:`dpctl.compiler` instead.
 
 :py:mod:`dpctl.program` provides a way to create a SYCL kernel
 from either an OpenCL* program source code represented as a string
@@ -20,20 +23,8 @@ execution via :py:meth:`dpctl.SyclQueue.submit`.
     :toctree: generated
     :nosignatures:
 
-    create_kernel_bundle_from_source
-    create_kernel_bundle_from_spirv
     create_program_from_source
     create_program_from_spirv
-
-.. autosummary::
-    :toctree: generated
-    :nosignatures:
-
-    SyclKernelBundle
     SyclKernel
-
-.. autosummary::
-    :toctree: generated
-    :nosignatures:
-
-    SyclKernelBundleCompilationError
+    SyclProgram
+    SyclProgramCompilationError

docs/doc_sources/api_reference/index.rst

@@ -9,7 +9,8 @@ The package ``dpctl`` provides
 * Python language bindings for the DPC++ runtime
    - :ref:`API objects <dpctl_pyapi>` in :py:mod:`dpctl` namespace
    - :ref:`API objects <dpctl_memory_pyapi>` in :py:mod:`dpctl.memory` namespace
-   - :ref:`API objects <dpctl_program_pyapi>` in :py:mod:`dpctl.program` namespace
+   - :ref:`API objects <dpctl_compiler_pyapi>` in :py:mod:`dpctl.compiler` namespace
+   - :ref:`API objects <dpctl_program_pyapi>` in :py:mod:`dpctl.program` namespace (deprecated, use :py:mod:`dpctl.compiler`)
    - :ref:`API objects <dpctl_utils_pyapi>` in :py:mod:`dpctl.utils` namespace
 * Python C API
    - :ref:`C API <dpctl_capi>` for working with Python classes defined in :mod:`dpctl`
@@ -31,6 +32,7 @@ The package ``dpctl`` provides
 
    dpctl/index
    dpctl/memory
+   dpctl/compiler
    dpctl/program
    dpctl/utils
    libsyclinterface/index

dpctl/CMakeLists.txt

@@ -203,6 +203,6 @@ build_dpctl_ext(${_trgt} ${_cy_file} "dpctl" SYCL)
 target_include_directories(${_trgt} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR})
 target_link_libraries(DpctlCAPI INTERFACE ${_trgt}_headers)
 
-add_subdirectory(program)
+add_subdirectory(compiler)
 add_subdirectory(memory)
 add_subdirectory(utils)

dpctl/__init__.py

@@ -125,6 +125,7 @@
 # add submodules
 __all__ += [
     "memory",
+    "compiler",
     "program",
     "utils",
 ]

dpctl/_sycl_queue.pxd

@@ -32,7 +32,7 @@ from ._backend cimport (
 from ._sycl_context cimport SyclContext
 from ._sycl_device cimport SyclDevice
 from ._sycl_event cimport SyclEvent
-from .program._program cimport SyclKernel
+from .compiler._compiler cimport SyclKernel
 
 
 cdef public api class _SyclQueue [

dpctl/_sycl_queue.pyx

@@ -383,7 +383,7 @@ kernel_arg_type = _kernel_arg_type()
 cdef class SyclKernelSubmitError(Exception):
     """
     A ``SyclKernelSubmitError`` exception is raised when
-    the provided :class:`.program.SyclKernel` could not be
+    the provided :class:`.compiler.SyclKernel` could not be
     submitted to the :class:`.SyclQueue`.
 
     """
@@ -1178,10 +1178,10 @@ cdef class SyclQueue(_SyclQueue):
         list dEvents=None
     ):
         """
-        Asynchronously submit :class:`dpctl.program.SyclKernel` for execution.
+        Asynchronously submit :class:`dpctl.compiler.SyclKernel` for execution.
 
         Args:
-            kernel (dpctl.program.SyclKernel):
+            kernel (dpctl.compiler.SyclKernel):
                 SYCL kernel object
             args (List[object]):
                 List of kernel arguments
@@ -1339,10 +1339,10 @@ cdef class SyclQueue(_SyclQueue):
         list dEvents=None
     ):
         """
-        Submit :class:`dpctl.program.SyclKernel` for execution.
+        Submit :class:`dpctl.compiler.SyclKernel` for execution.
 
         Args:
-            kernel (dpctl.program.SyclKernel):
+            kernel (dpctl.compiler.SyclKernel):
                 SYCL kernel object
             args (List[object]):
                 List of kernel arguments

dpctl/apis/include/dpctl4pybind11.hpp

@@ -227,7 +227,7 @@ class dpctl_capi
         this->Memory_GetNumBytes_ = Memory_GetNumBytes;
         this->Memory_Make_ = Memory_Make;
 
-        // dpctl.program API
+        // dpctl.compiler API
         this->SyclKernel_GetKernelRef_ = SyclKernel_GetKernelRef;
         this->SyclKernel_Make_ = SyclKernel_Make;
         this->SyclKernelBundle_GetKernelBundleRef_ =
@@ -450,7 +450,7 @@ template <> struct type_caster<sycl::event>
 };
 
 /* This type caster associates ``sycl::kernel`` C++ class with
- * :class:`dpctl.program.SyclKernel` for the purposes of generation of
+ * :class:`dpctl.compiler.SyclKernel` for the purposes of generation of
  * Python bindings by pybind11.
  */
 template <> struct type_caster<sycl::kernel>
@@ -469,7 +469,7 @@ template <> struct type_caster<sycl::kernel>
         }
         else {
             throw py::type_error("Input is of unexpected type, expected "
-                                 "dpctl.program.SyclKernel");
+                                 "dpctl.compiler.SyclKernel");
         }
     }
 
@@ -482,12 +482,12 @@ template <> struct type_caster<sycl::kernel>
         return handle(reinterpret_cast<PyObject *>(tmp));
     }
 
-    DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.program.SyclKernel"));
+    DPCTL_TYPE_CASTER(sycl::kernel, _("dpctl.compiler.SyclKernel"));
 };
 
 /* This type caster associates
  * ``sycl::kernel_bundle<sycl::bundle_state::executable>`` C++ class with
- * :class:`dpctl.program.SyclKernelBundle` for the purposes of generation of
+ * :class:`dpctl.compiler.SyclKernelBundle` for the purposes of generation of
  * Python bindings by pybind11.
  */
 template <>
@@ -511,7 +511,7 @@ struct type_caster<sycl::kernel_bundle<sycl::bundle_state::executable>>
         }
         else {
             throw py::type_error("Input is of unexpected type, expected "
-                                 "dpctl.program.SyclKernelBundle");
+                                 "dpctl.compiler.SyclKernelBundle");
         }
     }
 
@@ -526,7 +526,7 @@ struct type_caster<sycl::kernel_bundle<sycl::bundle_state::executable>>
     }
 
     DPCTL_TYPE_CASTER(sycl::kernel_bundle<sycl::bundle_state::executable>,
-                      _("dpctl.program.SyclKernelBundle"));
+                      _("dpctl.compiler.SyclKernelBundle"));
 };
 
 /* This type caster associates

dpctl/apis/include/dpctl_capi.h

@@ -45,8 +45,8 @@
 #include "dpctl/_sycl_queue_api.h"
 #include "dpctl/memory/_memory.h"
 #include "dpctl/memory/_memory_api.h"
-#include "dpctl/program/_program.h"
-#include "dpctl/program/_program_api.h"
+#include "dpctl/compiler/_compiler.h"
+#include "dpctl/compiler/_compiler_api.h"
 
 // clang-format on
 
@@ -66,6 +66,6 @@ static inline void import_dpctl(void)
     import_dpctl___sycl_event();
     import_dpctl___sycl_queue();
     import_dpctl__memory___memory();
-    import_dpctl__program___program();
+    import_dpctl__compiler___compiler();
     return;
 }

dpctl/program/CMakeLists.txt → dpctl/compiler/CMakeLists.txt

@@ -1,7 +1,6 @@
-
 file(GLOB _cython_sources *.pyx)
 foreach(_cy_file ${_cython_sources})
     get_filename_component(_trgt ${_cy_file} NAME_WLE)
-    build_dpctl_ext(${_trgt} ${_cy_file} "dpctl/program" RELATIVE_PATH "..")
+    build_dpctl_ext(${_trgt} ${_cy_file} "dpctl/compiler" RELATIVE_PATH "..")
     target_link_libraries(DpctlCAPI INTERFACE ${_trgt}_headers)
 endforeach()

dpctl/compiler/__init__.pxd

@@ -0,0 +1,32 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2025 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# distutils: language = c++
+# cython: language_level=3
+
+"""Declares the extension types and functions for the Cython API
+implemented in dpctl.compiler._compiler.pyx.
+"""
+
+
+from dpctl.compiler._compiler cimport (
+    SyclKernel,
+    SyclKernelBundle,
+    create_kernel_bundle_from_source,
+    create_kernel_bundle_from_spirv,
+    create_program_from_source,
+    create_program_from_spirv,
+)

dpctl/compiler/__init__.py

@@ -0,0 +1,45 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2025 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""
+**Data Parallel Control Compiler** provides a way to create a SYCL kernel
+from either an OpenCL program represented as a string or a SPIR-V binary
+file.
+
+"""
+
+from dpctl.compiler._compiler import (
+    SpecializationConstant,
+    SyclKernel,
+    SyclKernelBundle,
+    SyclKernelBundleCompilationError,
+    create_kernel_bundle_from_source,
+    create_kernel_bundle_from_spirv,
+)
+
+__all__ = [
+    "create_kernel_bundle_from_source",
+    "create_kernel_bundle_from_spirv",
+    "SyclKernel",
+    "SyclKernelBundle",
+    "SyclKernelBundleCompilationError",
+    "SpecializationConstant",
+]
+
+# add submodules
+__all__ += [
+    "utils",
+]

dpctl/program/_program.pxd → dpctl/compiler/_compiler.pxd

@@ -68,7 +68,3 @@ cpdef create_kernel_bundle_from_spirv (
     unicode copts=*,
     list specializations=*,
 )
-cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*)
-cpdef create_program_from_spirv (
-    SyclQueue q, const unsigned char[:] IL, unicode copts=*
-)