Merge pull request #199 from diptorupd/feature/program

Feature/program
IntelPython · Dec 4, 2020 · bbd586a · bbd586a
2 parents 45959f9 + fdfdcaf
commit bbd586a
Show file tree

Hide file tree

Showing 25 changed files with 868 additions and 695 deletions.
diff --git a/dpctl-capi/include/dpctl_sycl_program_interface.h b/dpctl-capi/include/dpctl_sycl_program_interface.h
@@ -41,11 +41,10 @@ DPCTL_C_EXTERN_C_BEGIN
  * @brief Create a Sycl program from an OpenCL SPIR-V binary file.
  *
  * Sycl 1.2 does not expose any method to create a sycl::program from a SPIR-V
- * IL file. To get around this limitation, we need to use the Sycl feature to
- * create an interoperability kernel from an OpenCL kernel. This function first
- * creates an OpenCL program and kernel from the SPIR-V binary and then using
- * the Sycl-OpenCL interoperability feature creates a Sycl kernel from the
- * OpenCL kernel.
+ * IL file. To get around this limitation, we first creare a SYCL
+ * interoperability program and then create a SYCL program from the
+ * interoperability program. Currently, interoperability programs can be created
+ * for OpenCL and Level-0 backends.
  *
  * The feature to create a Sycl kernel from a SPIR-V IL binary will be available
  * in Sycl 2.0 at which point this function may become deprecated.

diff --git a/dpctl-capi/source/dpctl_sycl_program_interface.cpp b/dpctl-capi/source/dpctl_sycl_program_interface.cpp
@@ -38,30 +38,21 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPCTLSyclContextRef)
 DEFINE_SIMPLE_CONVERSION_FUNCTIONS(program, DPCTLSyclProgramRef)
 DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef)
 
-} /* end of anonymous namespace */
-
 __dpctl_give DPCTLSyclProgramRef
-DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef,
-                                __dpctl_keep const void *IL,
-                                size_t length)
+createOpenCLInterOpProgram (const context &SyclCtx,
+                            __dpctl_keep const void *IL,
+                            size_t length)
 {
     cl_int err;
-    context *SyclCtx;
-    if(!CtxRef) {
-        // \todo handle error
-        return nullptr;
-    }
-
-    SyclCtx = unwrap(CtxRef);
-    auto CLCtx   = SyclCtx->get();
+    auto CLCtx   = SyclCtx.get();
     auto CLProgram = clCreateProgramWithIL(CLCtx, IL, length, &err);
     if (err) {
         // \todo: record the error string and any other information.
         std::cerr << "OpenCL program could not be created from the SPIR-V "
                      "binary. OpenCL Error " << err << ".\n";
         return nullptr;
     }
-    auto SyclDevices = SyclCtx->get_devices();
+    auto SyclDevices = SyclCtx.get_devices();
 
     // Get a list of CL Devices from the Sycl devices
     auto CLDevices = new cl_device_id[SyclDevices.size()];
@@ -83,18 +74,50 @@ DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef,
 
     // Create the Sycl program from OpenCL program
     try {
-        auto SyclProgram = new program(*SyclCtx, CLProgram);
+        auto SyclProgram = new program(SyclCtx, CLProgram);
         return wrap(SyclProgram);
-    } catch (invalid_object_error) {
+    } catch (invalid_object_error &e) {
         // \todo record error
+        std::cerr << e.what() << '\n';
         return nullptr;
     }
 }
 
+} /* end of anonymous namespace */
+
+__dpctl_give DPCTLSyclProgramRef
+DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef,
+                                 __dpctl_keep const void *IL,
+                                 size_t length)
+{
+    DPCTLSyclProgramRef Pref = nullptr;
+    context *SyclCtx = nullptr;
+    if(!CtxRef) {
+        // \todo handle error
+        return Pref;
+    }
+
+    SyclCtx = unwrap(CtxRef);
+    // get the backend type
+    auto BE = SyclCtx->get_platform().get_backend();
+    switch (BE)
+    {
+    case backend::opencl:
+        Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length);
+        break;
+    case backend::level_zero:
+        break;
+    default:
+        break;
+    }
+
+    return Pref;
+}
+
 __dpctl_give DPCTLSyclProgramRef
 DPCTLProgram_CreateFromOCLSource (__dpctl_keep const DPCTLSyclContextRef Ctx,
-                                 __dpctl_keep const char *Source,
-                                 __dpctl_keep const char *CompileOpts)
+                                  __dpctl_keep const char *Source,
+                                  __dpctl_keep const char *CompileOpts)
 {
     std::string compileOpts;
     context *SyclCtx = nullptr;
@@ -118,23 +141,43 @@ DPCTLProgram_CreateFromOCLSource (__dpctl_keep const DPCTLSyclContextRef Ctx,
         compileOpts = CompileOpts;
     }
 
-    try{
-        SyclProgram->build_with_source(source, compileOpts);
-        return wrap(SyclProgram);
-    } catch (compile_program_error) {
-        delete SyclProgram;
-        // \todo record error
+    // get the backend type
+    auto BE = SyclCtx->get_platform().get_backend();
+    switch (BE)
+    {
+    case backend::opencl:
+        try {
+            SyclProgram->build_with_source(source, compileOpts);
+            return wrap(SyclProgram);
+        } catch (compile_program_error &e) {
+            std::cerr << e.what() << '\n';
+            delete SyclProgram;
+            // \todo record error
+            return nullptr;
+        } catch (feature_not_supported &e) {
+            std::cerr << e.what() << '\n';
+            delete SyclProgram;
+            // \todo record error
+            return nullptr;
+        } catch (runtime_error &e) {
+            std::cerr << e.what() << '\n';
+            delete SyclProgram;
+            // \todo record error
+            return nullptr;
+        }
+        break;
+    case backend::level_zero:
+        std::cerr << "CreateFromSource is not supported in Level Zero.\n";
         return nullptr;
-    } catch (feature_not_supported) {
-        delete SyclProgram;
-        // \todo record error
+    default:
+        std::cerr << "CreateFromSource is not supported in unknown backend.\n";
         return nullptr;
     }
 }
 
 __dpctl_give DPCTLSyclKernelRef
 DPCTLProgram_GetKernel (__dpctl_keep DPCTLSyclProgramRef PRef,
-                       __dpctl_keep const char *KernelName)
+                        __dpctl_keep const char *KernelName)
 {
     if(!PRef) {
         // \todo record error
@@ -149,15 +192,16 @@ DPCTLProgram_GetKernel (__dpctl_keep DPCTLSyclProgramRef PRef,
     try {
         auto SyclKernel = new kernel(SyclProgram->get_kernel(name));
         return wrap(SyclKernel);
-    } catch (invalid_object_error) {
+    } catch (invalid_object_error &e) {
         // \todo record error
+        std::cerr << e.what() << '\n';
         return nullptr;
     }
 }
 
 bool
 DPCTLProgram_HasKernel (__dpctl_keep DPCTLSyclProgramRef PRef,
-                       __dpctl_keep const char *KernelName)
+                        __dpctl_keep const char *KernelName)
 {
     if(!PRef) {
         // \todo handle error
@@ -167,7 +211,8 @@ DPCTLProgram_HasKernel (__dpctl_keep DPCTLSyclProgramRef PRef,
     auto SyclProgram = unwrap(PRef);
     try {
         return SyclProgram->has_kernel(KernelName);
-    } catch (invalid_object_error) {
+    } catch (invalid_object_error &e) {
+        std::cerr << e.what() << '\n';
         return false;
     }
 }

diff --git a/dpctl/__init__.pxd b/dpctl/__init__.pxd
@@ -1,32 +1,31 @@
-##===-------------- __init__.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.
-##
-##===----------------------------------------------------------------------===##
+# ===------------ __init__.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 dpctl._sycl_core cimport *
 from dpctl._memory import *
-
diff --git a/dpctl/__init__.py b/dpctl/__init__.py
@@ -1,27 +1,27 @@
-##===----------------- __init__.py - 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 top-level dpctl module.
-##
-##===----------------------------------------------------------------------===##
+# ===---------------- __init__.py - 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
+# The top-level dpctl module.
+#
+# ===-----------------------------------------------------------------------===#
 """
     **Data Parallel Control (dpCtl)**
 

diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd
@@ -1,28 +1,28 @@
-##===------------- backend.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 defines the Cython extern types for the functions and opaque data
-## types defined by dpctl's C API.
-##
-##===----------------------------------------------------------------------===##
+# ===------------ backend.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 defines the Cython extern types for the functions and opaque data
+# types defined by dpctl's C API.
+#
+# ===-----------------------------------------------------------------------===#
 
 # distutils: language = c++
 # cython: language_level=3