Move SyclKernel, SyclProgram, program creation functions into a separate sub-module.

Author: Diptorup Deb

dpctl-capi/source/dpctl_sycl_program_interface.cpp

@@ -38,9 +38,9 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPCTLSyclContextRef)
 DEFINE_SIMPLE_CONVERSION_FUNCTIONS(program, DPCTLSyclProgramRef)
 DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef)
 
-__dppl_give DPCTLSyclProgramRef
+__dpctl_give DPCTLSyclProgramRef
 createOpenCLInterOpProgram (const context &SyclCtx,
-                            __dppl_keep const void *IL,
+                            __dpctl_keep const void *IL,
                             size_t length)
 {
     cl_int err;
@@ -85,12 +85,12 @@ createOpenCLInterOpProgram (const context &SyclCtx,
 
 } /* end of anonymous namespace */
 
-__dppl_give DPCTLSyclProgramRef
-DPCTLProgram_CreateFromOCLSpirv (__dppl_keep const DPPLSyclContextRef CtxRef,
-                                 __dppl_keep const void *IL,
+__dpctl_give DPCTLSyclProgramRef
+DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef,
+                                 __dpctl_keep const void *IL,
                                  size_t length)
 {
-    DPPLSyclProgramRef Pref = nullptr;
+    DPCTLSyclProgramRef Pref = nullptr;
     context *SyclCtx = nullptr;
     if(!CtxRef) {
         // \todo handle error
@@ -114,10 +114,10 @@ DPCTLProgram_CreateFromOCLSpirv (__dppl_keep const DPPLSyclContextRef CtxRef,
     return Pref;
 }
 
-__dppl_give DPCTLSyclProgramRef
-DPCTLProgram_CreateFromOCLSource (__dppl_keep const DPPLSyclContextRef Ctx,
-                                  __dppl_keep const char *Source,
-                                  __dppl_keep const char *CompileOpts)
+__dpctl_give DPCTLSyclProgramRef
+DPCTLProgram_CreateFromOCLSource (__dpctl_keep const DPCTLSyclContextRef Ctx,
+                                  __dpctl_keep const char *Source,
+                                  __dpctl_keep const char *CompileOpts)
 {
     std::string compileOpts;
     context *SyclCtx = nullptr;

dpctl/_sycl_core.pxd

@@ -1,33 +1,34 @@
-##===------------- sycl_core.pxd - dpctl module --------*- Cython -*-------===##
-##
-##                      Data Parallel Control (dpCtl)
-##
-## Copyright 2020 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.
-##
-##===----------------------------------------------------------------------===##
-##
-## \file
-## This file declares the extension types and functions for the Cython API
-## implemented in sycl_core.pyx.
-##
-##===----------------------------------------------------------------------===##
+# ===------------- sycl_core.pxd - dpctl module --------*- Cython -*--------===#
+#
+#                      Data Parallel Control (dpCtl)
+#
+# Copyright 2020 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.
+#
+# ===-----------------------------------------------------------------------===#
+#
+# \file
+# This file declares the extension types and functions for the Cython API
+# implemented in sycl_core.pyx.
+#
+# ===-----------------------------------------------------------------------===#
 
 # distutils: language = c++
 # cython: language_level=3
 
 from ._backend cimport *
+from .program._program cimport SyclKernel
 from libc.stdint cimport uint32_t
 
 
@@ -85,34 +86,6 @@ cdef class SyclEvent:
     cpdef void wait (self)
 
 
-cdef class SyclKernel:
-    ''' Wraps a sycl::kernel object created from an OpenCL interoperability
-        kernel.
-    '''
-    cdef DPCTLSyclKernelRef _kernel_ref
-    cdef const char *_function_name
-    cdef DPCTLSyclKernelRef get_kernel_ref (self)
-
-    @staticmethod
-    cdef SyclKernel _create (DPCTLSyclKernelRef kref)
-
-
-cdef class SyclProgram:
-    ''' Wraps a sycl::program object created from an OpenCL interoperability
-        program.
-
-        SyclProgram exposes the C API from dpctl_sycl_program_interface.h. A
-        SyclProgram can be created from either a source string or a SPIR-V
-        binary file.
-    '''
-    cdef DPCTLSyclProgramRef _program_ref
-
-    @staticmethod
-    cdef  SyclProgram _create (DPCTLSyclProgramRef pref)
-    cdef  DPCTLSyclProgramRef get_program_ref (self)
-    cpdef SyclKernel get_sycl_kernel(self, str kernel_name)
-
-
 cdef class SyclQueue:
     ''' Wrapper class for a Sycl queue.
     '''

dpctl/_sycl_core.pyx

@@ -1,27 +1,27 @@
-##===------------- sycl_core.pyx - dpctl module -------*- Cython -*--------===##
-##
-##                      Data Parallel Control (dpCtl)
-##
-## Copyright 2020 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.
-##
-##===----------------------------------------------------------------------===##
-##
-## \file
-## This file implements a sub-set of Sycl's interface using dpctl's CAPI.
-##
-##===----------------------------------------------------------------------===##
+# ===------------ sycl_core.pyx - dpctl module -------*- Cython -*----------===#
+#
+#                      Data Parallel Control (dpCtl)
+#
+# Copyright 2020 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.
+#
+# ===-----------------------------------------------------------------------===#
+#
+# \file
+# This file implements a sub-set of Sycl's interface using dpctl's CAPI.
+#
+# ===-----------------------------------------------------------------------===#
 
 # distutils: language = c++
 # cython: language_level=3
@@ -46,16 +46,12 @@ __all__ = [
     "has_sycl_platforms",
     "set_default_queue",
     "is_in_device_context",
-    "create_program_from_source",
-    "create_program_from_spirv",
     "device_type",
     "backend_type",
     "device_context",
     "SyclContext",
     "SyclDevice",
     "SyclEvent",
-    "SyclKernel",
-    "SyclProgram",
     "SyclQueue"
 ]
 
@@ -92,12 +88,6 @@ cdef class UnsupportedDeviceError (Exception):
     """
     pass
 
-cdef class SyclProgramCompilationError (Exception):
-    """This exception is raised when a sycl program could not be built from
-       either a spirv binary file or a string source.
-    """
-    pass
-
 cdef class SyclKernelSubmitError (Exception):
     """This exception is raised when a sycl program could not be built from
        either a spirv binary file or a string source.
@@ -306,88 +296,6 @@ cdef class SyclEvent:
         """
         return int(<size_t>self._event_ref)
 
-
-cdef class SyclKernel:
-    """ Wraps a sycl::kernel object created from an OpenCL interoperability
-        kernel.
-    """
-
-    @staticmethod
-    cdef SyclKernel _create (DPCTLSyclKernelRef kref):
-        cdef SyclKernel ret = SyclKernel.__new__(SyclKernel)
-        ret._kernel_ref = kref
-        ret._function_name = DPCTLKernel_GetFunctionName(kref)
-        return ret
-
-    def __dealloc__ (self):
-        DPCTLKernel_Delete(self._kernel_ref)
-        DPCTLCString_Delete(self._function_name)
-
-    def get_function_name (self):
-        """ Returns the name of the Kernel function.
-        """
-        return self._function_name.decode()
-
-    def get_num_args (self):
-        """ Returns the number of arguments for this kernel function.
-        """
-        return DPCTLKernel_GetNumArgs(self._kernel_ref)
-
-    cdef DPCTLSyclKernelRef get_kernel_ref (self):
-        """ Returns the DPCTLSyclKernelRef pointer for this SyclKernel.
-        """
-        return self._kernel_ref
-
-    def addressof_ref (self):
-        """ Returns the address of the C API DPCTLSyclKernelRef pointer
-        as a long.
-
-        Returns:
-            The address of the DPCTLSyclKernelRef object used to create this
-            SyclKernel cast to a long.
-        """
-        return int(<size_t>self._kernel_ref)
-
-cdef class SyclProgram:
-    """ Wraps a sycl::program object created from an OpenCL interoperability
-        program.
-
-        SyclProgram exposes the C API from dpctl_sycl_program_interface.h. A
-        SyclProgram can be created from either a source string or a SPIR-V
-        binary file.
-    """
-
-    @staticmethod
-    cdef SyclProgram _create (DPCTLSyclProgramRef pref):
-        cdef SyclProgram ret = SyclProgram.__new__(SyclProgram)
-        ret._program_ref = pref
-        return ret
-
-    def __dealloc__ (self):
-        DPCTLProgram_Delete(self._program_ref)
-
-    cdef DPCTLSyclProgramRef get_program_ref (self):
-        return self._program_ref
-
-    cpdef SyclKernel get_sycl_kernel(self, str kernel_name):
-        name = kernel_name.encode('utf8')
-        return SyclKernel._create(DPCTLProgram_GetKernel(self._program_ref,
-                                                        name))
-
-    def has_sycl_kernel(self, str kernel_name):
-        name = kernel_name.encode('utf8')
-        return DPCTLProgram_HasKernel(self._program_ref, name)
-
-    def addressof_ref (self):
-        """Returns the address of the C API DPCTLSyclProgramRef pointer
-        as a long.
-
-        Returns:
-            The address of the DPCTLSyclProgramRef object used to create this
-            SyclProgram cast to a long.
-        """
-        return int(<size_t>self._program_ref)
-
 import ctypes
 
 cdef class SyclQueue:
@@ -894,80 +802,6 @@ cpdef get_current_backend():
     """
     return _mgr.get_current_backend()
 
-def create_program_from_source (SyclQueue q, unicode source, unicode copts=""):
-    """
-        Creates a Sycl interoperability program from an OpenCL source string.
-
-        We use the DPCTLProgram_CreateFromOCLSource() C API function to create
-        a Sycl progrma from an OpenCL source program that can contain multiple
-        kernels.
-
-        Parameters:
-            q (SyclQueue)   : The :class:`SyclQueue` for which the
-                              :class:`SyclProgram` is going to be built.
-            source (unicode): Source string for an OpenCL program.
-            copts (unicode) : Optional compilation flags that will be used
-                              when compiling the program.
-
-        Returns:
-            program (SyclProgram): A :class:`SyclProgram` object wrapping the  sycl::program returned by the C API.
-    """
-
-    BE = q.get_sycl_backend()
-    if BE != backend_type.opencl:
-        raise ValueError(
-            "Cannot create program for a ", BE, "type backend. Currently only "
-            "OpenCL devices are supported for program creations."
-        )
-
-    cdef DPCTLSyclProgramRef Pref
-    cdef bytes bSrc = source.encode('utf8')
-    cdef bytes bCOpts = copts.encode('utf8')
-    cdef const char *Src = <const char*>bSrc
-    cdef const char *COpts = <const char*>bCOpts
-    cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
-    Pref = DPCTLProgram_CreateFromOCLSource(CRef, Src, COpts)
-
-    if Pref is NULL:
-        raise SyclProgramCompilationError()
-
-    return SyclProgram._create(Pref)
-
-cimport cython.array
-
-def create_program_from_spirv (SyclQueue q, const unsigned char[:] IL):
-    """
-        Creates a Sycl interoperability program from an SPIR-V binary.
-
-        We use the DPCTLProgram_CreateFromOCLSpirv() C API function to create
-        a Sycl progrma from an compiled SPIR-V binary file.
-
-        Parameters:
-            q (SyclQueue): The :class:`SyclQueue` for which the
-                           :class:`SyclProgram` is going to be built.
-            IL (const char[:]) : SPIR-V binary IL file for an OpenCL program.
-
-        Returns:
-            program (SyclProgram): A :class:`SyclProgram` object wrapping the  sycl::program returned by the C API.
-    """
-    BE = q.get_sycl_backend()
-    if BE != backend_type.opencl:
-        raise ValueError(
-            "Cannot create program for a ", BE, "type backend. Currently only "
-            "OpenCL devices are supported for program creations."
-        )
-
-    cdef DPCTLSyclProgramRef Pref
-    cdef const unsigned char *dIL = &IL[0]
-    cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
-    cdef size_t length = IL.shape[0]
-    Pref = DPCTLProgram_CreateFromOCLSpirv(CRef, <const void*>dIL, length)
-    if Pref is NULL:
-        raise SyclProgramCompilationError()
-
-    return SyclProgram._create(Pref)
-
-
 from contextlib import contextmanager
 
 @contextmanager