diff --git a/MANIFEST.in b/MANIFEST.in index bbf310133c..43d8bddba4 100644 --- a/MANIFEST.in +++ b/MANIFEST.in @@ -2,4 +2,5 @@ include versioneer.py recursive-include dpctl/include *.h *.hpp include dpctl/*.pxd include dpctl/*DPPL*Interface.* +include dpctl/tests/input_files/* global-exclude *.cpp diff --git a/backends/CMakeLists.txt b/backends/CMakeLists.txt index c92faed9b8..373f0c8f2b 100644 --- a/backends/CMakeLists.txt +++ b/backends/CMakeLists.txt @@ -98,7 +98,10 @@ add_library( SHARED source/dppl_sycl_context_interface.cpp source/dppl_sycl_device_interface.cpp + source/dppl_sycl_event_interface.cpp + source/dppl_sycl_kernel_interface.cpp source/dppl_sycl_platform_interface.cpp + source/dppl_sycl_program_interface.cpp source/dppl_sycl_queue_interface.cpp source/dppl_sycl_queue_manager.cpp source/dppl_sycl_usm_interface.cpp @@ -146,13 +149,12 @@ if(WIN32) ) target_link_libraries( DPPLSyclInterface - PRIVATE - ${DPCPP_ROOT}/lib/sycl.lib + PRIVATE ${DPCPP_ROOT}/lib/sycl.lib + PRIVATE ${DPCPP_ROOT}/lib/OpenCL.lib ) target_link_libraries( DPPLOpenCLInterface - PRIVATE - ${DPCPP_ROOT}/lib/OpenCL.lib + PRIVATE ${DPCPP_ROOT}/lib/OpenCL.lib ) endif() diff --git a/backends/dbg_build.sh b/backends/dbg_build.sh index 6fba1486c9..2ae201b8cb 100755 --- a/backends/dbg_build.sh +++ b/backends/dbg_build.sh @@ -5,6 +5,7 @@ mkdir build pushd build INSTALL_PREFIX=`pwd`/../install +rm -rf ${INSTALL_PREFIX} export ONEAPI_ROOT=/opt/intel/oneapi DPCPP_ROOT=${ONEAPI_ROOT}/compiler/latest/linux PYTHON_INC=`python -c "import distutils.sysconfig; \ @@ -24,9 +25,6 @@ cmake \ -DGTEST_LIB_DIR=${CONDA_PREFIX}/lib \ .. -make V=1 -n -j 4 -make check -make install - +make V=1 -n -j 4 && make check && make install popd diff --git a/backends/include/dppl_sycl_event_interface.h b/backends/include/dppl_sycl_event_interface.h new file mode 100644 index 0000000000..c97eaf08f3 --- /dev/null +++ b/backends/include/dppl_sycl_event_interface.h @@ -0,0 +1,55 @@ +//===--- dppl_sycl_event_interface.h - DPPL-SYCL interface ---*---C++ -*---===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 header declares a C API to a sub-set of the sycl::event interface. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +#include "dppl_data_types.h" +#include "dppl_sycl_types.h" +#include "Support/DllExport.h" +#include "Support/ExternC.h" +#include "Support/MemOwnershipAttrs.h" + + +DPPL_C_EXTERN_C_BEGIN + +/*! + * @brief C-API wrapper for sycl::event.wait. + * + * @param ERef An opaque DPPLSyclEventRef pointer on which to wait. + */ +DPPL_API +void DPPLEvent_Wait (__dppl_keep DPPLSyclEventRef ERef); + +/*! + * @brief Deletes the DPPLSyclEventRef after casting it to a sycl::event. + * + * @param ERef An opaque DPPLSyclEventRef pointer that would be + * freed. + */ +DPPL_API +void +DPPLEvent_Delete (__dppl_take DPPLSyclEventRef ERef); + +DPPL_C_EXTERN_C_END diff --git a/backends/include/dppl_sycl_kernel_interface.h b/backends/include/dppl_sycl_kernel_interface.h new file mode 100644 index 0000000000..64d2f97a30 --- /dev/null +++ b/backends/include/dppl_sycl_kernel_interface.h @@ -0,0 +1,74 @@ +//===---- dppl_sycl_kernel_interface.h - DPPL-SYCL interface --*--C++ --*--===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 header declares a C API to create Sycl kernels from OpenCL kernels. In +/// future, API to create interoperability kernels from other languages such as +/// Level-0 driver API may be added here. +/// +/// \todo Investigate what we should do when we add support for Level-0 API. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +#include "dppl_data_types.h" +#include "dppl_sycl_types.h" +#include "Support/DllExport.h" +#include "Support/ExternC.h" +#include "Support/MemOwnershipAttrs.h" + +DPPL_C_EXTERN_C_BEGIN + +/*! + * @brief Returns a C string for the kernel name. + * + * @param KRef DPPLSyclKernelRef pointer to an OpenCL + * interoperability kernel. + * @return If a kernel name exists then returns it as a C string, else + * returns a nullptr. + */ +DPPL_API +__dppl_give const char* +DPPLKernel_GetFunctionName (__dppl_keep const DPPLSyclKernelRef KRef); + +/*! + * @brief Returns the number of arguments for the OpenCL kernel. + * + * @param KRef DPPLSyclKernelRef pointer to an OpenCL + * interoperability kernel. + * @return Returns the number of arguments for the OpenCL interoperability + * kernel. + */ +DPPL_API +size_t +DPPLKernel_GetNumArgs (__dppl_keep const DPPLSyclKernelRef KRef); + +/*! + * @brief Deletes the DPPLSyclKernelRef after casting it to a sycl::kernel. + * + * @param KRef DPPLSyclKernelRef pointer to an OpenCL + * interoperability kernel. + */ +DPPL_API +void +DPPLKernel_Delete (__dppl_take DPPLSyclKernelRef KRef); + +DPPL_C_EXTERN_C_END diff --git a/backends/include/dppl_sycl_program_interface.h b/backends/include/dppl_sycl_program_interface.h new file mode 100644 index 0000000000..e14f45b9f5 --- /dev/null +++ b/backends/include/dppl_sycl_program_interface.h @@ -0,0 +1,114 @@ +//===---- dppl_sycl_program_interface.h - DPPL-SYCL interface --*--C++ --*--===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 header declares a C API to create Sycl program an interoperability +/// program defined in OpenCL. In future, API to create interoperability +/// kernels from other languages such as Level-0 driver API may be added here. +/// +/// \todo Investigate what we should do when we add support for Level-0 API. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +#include "dppl_data_types.h" +#include "dppl_sycl_types.h" +#include "Support/DllExport.h" +#include "Support/ExternC.h" +#include "Support/MemOwnershipAttrs.h" + +DPPL_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. + * + * 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. + * + * @param Ctx An opaque pointer to a sycl::context + * @param IL SPIR-V binary + * @return A new SyclProgramRef pointer if the program creation succeeded, + * else returns NULL. + */ +DPPL_API +__dppl_give DPPLSyclProgramRef +DPPLProgram_CreateFromOCLSpirv (__dppl_keep const DPPLSyclContextRef Ctx, + __dppl_keep const void *IL, + size_t Length); + +/*! + * @brief Create a Sycl program from an OpenCL kernel source string. + * + * @param Ctx An opaque pointer to a sycl::context + * @param Source OpenCL source string + * @param CompileOptions Extra compiler flags (refer Sycl spec.) + * @return A new SyclProgramRef pointer if the program creation succeeded, + * else returns NULL. + */ +DPPL_API +__dppl_give DPPLSyclProgramRef +DPPLProgram_CreateFromOCLSource (__dppl_keep const DPPLSyclContextRef Ctx, + __dppl_keep const char *Source, + __dppl_keep const char *CompileOpts = nullptr); + +/*! + * @brief Returns the SyclKernel with given name from the program, if not found + * then return NULL. + * + * @param PRef Opaque pointer to a sycl::program + * @param KernelName Name of kernel + * @return A SyclKernel reference if the kernel exists, else NULL + */ +DPPL_API +__dppl_give DPPLSyclKernelRef +DPPLProgram_GetKernel (__dppl_keep DPPLSyclProgramRef PRef, + __dppl_keep const char *KernelName); + +/*! + * @brief Return True if a SyclKernel with given name exists in the program, if + * not found then returns False. + * + * @param PRef Opaque pointer to a sycl::program + * @param KernelName Name of kernel + * @return True if the kernel exists, else False + */ +DPPL_API +bool +DPPLProgram_HasKernel (__dppl_keep DPPLSyclProgramRef PRef, + __dppl_keep const char *KernelName); + +/*! + * @brief Frees the DPPLSyclProgramRef pointer. + * + * @param PRef Opaque pointer to a sycl::program + */ +DPPL_API +void +DPPLProgram_Delete (__dppl_take DPPLSyclProgramRef PRef); + +DPPL_C_EXTERN_C_END diff --git a/backends/include/dppl_sycl_queue_interface.h b/backends/include/dppl_sycl_queue_interface.h index 0aca3512be..cc12c9af77 100644 --- a/backends/include/dppl_sycl_queue_interface.h +++ b/backends/include/dppl_sycl_queue_interface.h @@ -35,6 +35,39 @@ DPPL_C_EXTERN_C_BEGIN +/*! + * @brief Supported types for kernel arguments to be passed to a Sycl kernel. + * + * \todo Add support for sycl::buffer + * + */ +typedef enum +{ + DPPL_CHAR, + DPPL_SIGNED_CHAR, + DPPL_UNSIGNED_CHAR, + DPPL_SHORT, + DPPL_INT, + DPPL_UNSIGNED_INT, + DPPL_LONG, + DPPL_UNSIGNED_LONG, + DPPL_LONG_LONG, + DPPL_UNSIGNED_LONG_LONG, + DPPL_SIZE_T, + DPPL_FLOAT, + DPPL_DOUBLE, + DPPL_LONG_DOUBLE, + DPPL_VOID_PTR +} DPPLKernelArgType; + +/*! + * @brief Delete the pointer after casting it to sycl::queue. + * + * @param QRef A DPPLSyclQueueRef pointer that gets deleted. + */ +DPPL_API +void DPPLQueue_Delete (__dppl_take DPPLSyclQueueRef QRef); + /*! * @brief Returns the Sycl context for the queue. * @@ -56,15 +89,109 @@ __dppl_give DPPLSyclDeviceRef DPPLQueue_GetDevice (__dppl_keep const DPPLSyclQueueRef QRef); /*! - * @brief Delete the pointer after casting it to sycl::queue. + * @brief Submits the kernel to the specified queue with the provided range + * argument. * - * @param QRef A DPPLSyclQueueRef pointer that gets deleted. + * A wrapper over sycl::queue.submit(). The function takes an interoperability + * kernel, the kernel arguments, and a Sycl queue as input. The kernel is + * submitted as parallel_for(range, *unwrap(KRef)). + * + * \todo sycl::buffer arguments are not supported yet. + * \todo Add support for id WorkItemOffset + * + * @param KRef Opaque pointer to an OpenCL interoperability kernel + * wrapped inside a sycl::kernel. + * @param QRef Opaque pointer to the sycl::queue where the kernel + * will be enqueued. + * @param Args An array of void* pointers that represent the + * kernel arguments for the kernel. + * @param ArgTypes An array of DPPLKernelArgType enum values that + * represent the type of each kernel argument. + * @param NArgs Size of Args and ArgTypes. + * @param Range Defines the overall dimension of the dispatch for + * the kernel. The array can have up to three + * dimensions. + * @param NRange Size of the gRange array. + * @param DepEvents List of dependent DPPLSyclEventRef objects (events) + * for the kernel. We call sycl::handler.depends_on for + * each of the provided events. + * @param NDepEvents Size of the DepEvents list. + * @return An opaque pointer to the sycl::event returned by the + * sycl::queue.submit() function. */ DPPL_API -void DPPLQueue_Delete (__dppl_take DPPLSyclQueueRef QRef); +DPPLSyclEventRef +DPPLQueue_SubmitRange (__dppl_keep const DPPLSyclKernelRef KRef, + __dppl_keep const DPPLSyclQueueRef QRef, + __dppl_keep void **Args, + __dppl_keep const DPPLKernelArgType *ArgTypes, + size_t NArgs, + __dppl_keep const size_t Range[3], + size_t NRange, + __dppl_keep const DPPLSyclEventRef *DepEvents, + size_t NDepEvents); + +/*! + * @brief Submits the kernel to the specified queue with the provided nd_range + * argument. + * + * A wrapper over sycl::queue.submit(). The function takes an interoperability + * kernel, the kernel arguments, and a Sycl queue as input. The kernel is + * submitted as parallel_for(nd_range, *unwrap(KRef)). + * + * \todo sycl::buffer arguments are not supported yet. + * \todo Add support for id WorkItemOffset + * + * @param KRef Opaque pointer to an OpenCL interoperability kernel + * wrapped inside a sycl::kernel. + * @param QRef Opaque pointer to the sycl::queue where the kernel + * will be enqueued. + * @param Args An array of void* pointers that represent the + * kernel arguments for the kernel. + * @param ArgTypes An array of DPPLKernelArgType enum values that + * represent the type of each kernel argument. + * @param NArgs Size of Args. + * @param gRange Defines the overall dimension of the dispatch for + * the kernel. The array can have up to three + * dimensions. + * @param lRange Defines the iteration domain of a single work-group + * in a parallel dispatch. The array can have up to + * three dimensions. + * @param NDims The number of dimensions for both local and global + * ranges. + * @param DepEvents List of dependent DPPLSyclEventRef objects (events) + * for the kernel. We call sycl::handler.depends_on for + * each of the provided events. + * @param NDepEvents Size of the DepEvents list. + * @return An opaque pointer to the sycl::event returned by the + * sycl::queue.submit() function. + */ +DPPL_API +DPPLSyclEventRef +DPPLQueue_SubmitNDRange(__dppl_keep const DPPLSyclKernelRef KRef, + __dppl_keep const DPPLSyclQueueRef QRef, + __dppl_keep void **Args, + __dppl_keep const DPPLKernelArgType *ArgTypes, + size_t NArgs, + __dppl_keep const size_t gRange[3], + __dppl_keep const size_t lRange[3], + size_t NDims, + __dppl_keep const DPPLSyclEventRef *DepEvents, + size_t NDepEvents); + +/*! + * @brief Calls the sycl::queue.submit function to do a blocking wait on all + * enqueued tasks in the queue. + * + * @param QRef Opaque pointer to a sycl::queue. + */ +DPPL_API +void +DPPLQueue_Wait (__dppl_keep const DPPLSyclQueueRef QRef); /*! - * @brief C-API wrapper for sycl::queue::memcpy. It waits an event. + * @brief C-API wrapper for sycl::queue::memcpy, the function waits on an event + * till the memcpy operation completes. * * @param QRef An opaque pointer to the sycl queue. * @param Dest An USM pointer to the destination memory. diff --git a/backends/include/dppl_sycl_types.h b/backends/include/dppl_sycl_types.h index 7d265805bd..e3ccf3d0a8 100644 --- a/backends/include/dppl_sycl_types.h +++ b/backends/include/dppl_sycl_types.h @@ -26,36 +26,48 @@ #pragma once /*! - * @brief + * @brief Opaque pointer to a sycl::context * */ typedef struct DPPLOpaqueSyclContext *DPPLSyclContextRef; /*! - * @brief + * @brief Opaque pointer to a sycl::device * */ typedef struct DPPLOpaqueSyclDevice *DPPLSyclDeviceRef; /*! - * @brief + * @brief Opaque pointer to a sycl::event + * + */ +typedef struct DPPLOpaqueSyclEvent *DPPLSyclEventRef; + +/*! + * @brief Opaque pointer to a sycl::kernel + * + */ +typedef struct DPPLOpaqueSyclKernel *DPPLSyclKernelRef; + +/*! + * @brief Opaque pointer to a sycl::platform * */ typedef struct DPPLOpaqueSyclPlatform *DPPLSyclPlatformRef; +/*! + * @brief Opaque pointer to a sycl::program + * + */ +typedef struct DPPLOpaqueSyclProgram *DPPLSyclProgramRef; + /*! - * @brief Used to pass a sycl::queue opaquely through DPPL interfaces. + * @brief Opaque pointer to a sycl::queue * * @see sycl::queue */ typedef struct DPPLOpaqueSyclQueue *DPPLSyclQueueRef; -/*! - * @brief Used to pass a sycl::program opaquely through DPPL interfaces. - * - */ -typedef struct DPPLOpaqueSyclProgram *DPPLSyclProgramRef; - /*! * @brief Used to pass a sycl::usm memory opaquely through DPPL interfaces. * diff --git a/backends/source/dppl_sycl_context_interface.cpp b/backends/source/dppl_sycl_context_interface.cpp index fc3aa26983..4d97b2bb6c 100644 --- a/backends/source/dppl_sycl_context_interface.cpp +++ b/backends/source/dppl_sycl_context_interface.cpp @@ -32,26 +32,16 @@ using namespace cl::sycl; namespace { - // Create wrappers for C Binding types (see CBindingWrapping.h). - DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPPLSyclContextRef) +// Create wrappers for C Binding types (see CBindingWrapping.h). +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPPLSyclContextRef) } /* end of anonymous namespace */ -/*! - * @brief - * - * @param CtxtRef My Param doc - * @return {return} My Param doc - */ + bool DPPLContext_IsHost (__dppl_keep const DPPLSyclContextRef CtxRef) { return unwrap(CtxRef)->is_host(); } -/*! - * @brief - * - * @param CtxtRef My Param doc - */ void DPPLContext_Delete (__dppl_take DPPLSyclContextRef CtxRef) { delete unwrap(CtxRef); diff --git a/backends/source/dppl_sycl_device_interface.cpp b/backends/source/dppl_sycl_device_interface.cpp index 1d9d9444c4..68e50ab9b0 100644 --- a/backends/source/dppl_sycl_device_interface.cpp +++ b/backends/source/dppl_sycl_device_interface.cpp @@ -35,7 +35,7 @@ using namespace cl::sycl; namespace { // Create wrappers for C Binding types (see CBindingWrapping.h). - DEFINE_SIMPLE_CONVERSION_FUNCTIONS(device, DPPLSyclDeviceRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(device, DPPLSyclDeviceRef) /*! * @brief Helper function to print the metadata for a sycl::device. @@ -54,6 +54,23 @@ void dump_device_info (const device & Device) << Device.get_info() << '\n'; ss << std::setw(4) << " " << std::left << std::setw(16) << "Profile" << Device.get_info() << '\n'; + ss << std::setw(4) << " " << std::left << std::setw(16) << "Device type"; + + try { + if (Device.has(aspect::accelerator)) + ss << "accelerator" << '\n'; + else if (Device.has(aspect::cpu)) + ss << "cpu" << '\n'; + else if (Device.has(aspect::custom)) + ss << "custom" << '\n'; + else if (Device.has(aspect::gpu)) + ss << "gpu" << '\n'; + else if (Device.has(aspect::host)) + ss << "host" << '\n'; + } catch (runtime_error re) { + // \todo handle errors + ss << "unknown\n"; + } std::cout << ss.str(); } diff --git a/backends/source/dppl_sycl_event_interface.cpp b/backends/source/dppl_sycl_event_interface.cpp new file mode 100644 index 0000000000..69a739483f --- /dev/null +++ b/backends/source/dppl_sycl_event_interface.cpp @@ -0,0 +1,52 @@ +//===--- dppl_sycl_event_interface.cpp - DPPL-SYCL interface --*- C++ -*---===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 the data types and functions declared in +/// dppl_sycl_event_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dppl_sycl_event_interface.h" +#include "Support/CBindingWrapping.h" + +#include /* SYCL headers */ + +using namespace cl::sycl; + +namespace +{ +// Create wrappers for C Binding types (see CBindingWrapping.h) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(event, DPPLSyclEventRef) +} /* end of anonymous namespace */ + + +void DPPLEvent_Wait (__dppl_keep DPPLSyclEventRef ERef) +{ + // \todo How to handle errors? E.g. when ERef is null or not a valid event. + auto SyclEvent = unwrap(ERef); + SyclEvent->wait(); +} + +void +DPPLEvent_Delete (__dppl_take DPPLSyclEventRef ERef) +{ + delete unwrap(ERef); +} diff --git a/backends/source/dppl_sycl_kernel_interface.cpp b/backends/source/dppl_sycl_kernel_interface.cpp new file mode 100644 index 0000000000..56fa9380cb --- /dev/null +++ b/backends/source/dppl_sycl_kernel_interface.cpp @@ -0,0 +1,75 @@ +//===--- dppl_sycl_kernel_interface.cpp - DPPL-SYCL interface --*-- C++ -*-===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 the functions declared in +/// dppl_sycl_kernel_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dppl_sycl_kernel_interface.h" +#include "Support/CBindingWrapping.h" + +#include /* Sycl headers */ + +using namespace cl::sycl; + +namespace +{ + +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPPLSyclKernelRef) + +} /* end of anonymous namespace */ + +__dppl_give const char* +DPPLKernel_GetFunctionName (__dppl_keep const DPPLSyclKernelRef Kernel) +{ + if(!Kernel) { + // \todo record error + return nullptr; + } + + auto SyclKernel = unwrap(Kernel); + auto kernel_name = SyclKernel->get_info(); + if(kernel_name.empty()) + return nullptr; + auto cstr_name = new char [kernel_name.length()+1]; + std::strcpy (cstr_name, kernel_name.c_str()); + return cstr_name; +} + +size_t +DPPLKernel_GetNumArgs (__dppl_keep const DPPLSyclKernelRef Kernel) +{ + if(!Kernel) { + // \todo record error + return -1; + } + + auto SyclKernel = unwrap(Kernel); + auto num_args = SyclKernel->get_info(); + return (size_t)num_args; +} + +void +DPPLKernel_Delete (__dppl_take DPPLSyclKernelRef Kernel) +{ + delete unwrap(Kernel); +} diff --git a/backends/source/dppl_sycl_platform_interface.cpp b/backends/source/dppl_sycl_platform_interface.cpp index 271a64b31a..09ec3cbb5e 100644 --- a/backends/source/dppl_sycl_platform_interface.cpp +++ b/backends/source/dppl_sycl_platform_interface.cpp @@ -38,31 +38,78 @@ using namespace cl::sycl; * found on the system: * - info::platform::name * - info::platform::version + * - info::platform::vendor * - info::platform::profile + * - backend (opencl, cuda, level-zero, host) + * - number of devices on the platform * + * Additionally, for each device we print out: + * - info::device::name + * - info::device::driver_version + * - type of the device based on the aspects cpu, gpu, accelerator. */ void DPPLPlatform_DumpInfo () { - size_t i = 0; + size_t i = 0; - // Print out the info for each platform - auto platforms = platform::get_platforms(); - for (auto &p : platforms) { - std::cout << "---Platform " << i << '\n'; - std::stringstream ss; + // Print out the info for each platform + auto platforms = platform::get_platforms(); + for (auto &p : platforms) { + std::cout << "---Platform " << i << '\n'; + std::stringstream ss; - ss << std::setw(4) << " " << std::left << std::setw(12) << "Name" - << p.get_info() << '\n'; - ss << std::setw(4) << " " << std::left << std::setw(12) << "Version" - << p.get_info() << '\n'; - ss << std::setw(4) << " " << std::left << std::setw(12) << "Vendor" - << p.get_info() << '\n'; - ss << std::setw(4) << " " << std::left << std::setw(12) << "Profile" - << p.get_info() << '\n'; + auto vendor = p.get_info(); + if (vendor.empty()) + vendor = "unknown"; - std::cout << ss.str(); - ++i; - } + ss << std::setw(4) << " " << std::left << std::setw(12) << "Name" + << p.get_info() << '\n'; + ss << std::setw(4) << " " << std::left << std::setw(12) << "Version" + << p.get_info() << '\n'; + ss << std::setw(4) << " " << std::left << std::setw(12) << "Vendor" + << vendor << '\n'; + ss << std::setw(4) << " " << std::left << std::setw(12) << "Profile" + << p.get_info() << '\n'; + ss << std::setw(4) << " " << std::left << std::setw(12) << "Backend"; + p.is_host() ? (ss << "unknown") : (ss << p.get_backend()); + ss << '\n'; + + // Get number of devices on the platform + auto devices = p.get_devices(); + + ss << std::setw(4) << " " << std::left << std::setw(12) << "Devices" + << devices.size() << '\n'; + // Print some of the device information + for (auto dn = 0ul; dn < devices.size(); ++dn) { + ss << std::setw(4) << "---Device " << dn << '\n'; + ss << std::setw(8) << " " << std::left << std::setw(20) + << "Name" << devices[dn].get_info() << '\n'; + ss << std::setw(8) << " " << std::left << std::setw(20) + << "Driver version" + << devices[dn].get_info() << '\n'; + ss << std::setw(8) << " " << std::left << std::setw(20) + << "Device type"; + + try { + if (devices[dn].has(aspect::accelerator)) + ss << "accelerator" << '\n'; + else if (devices[dn].has(aspect::cpu)) + ss << "cpu" << '\n'; + else if (devices[dn].has(aspect::custom)) + ss << "custom" << '\n'; + else if (devices[dn].has(aspect::gpu)) + ss << "gpu" << '\n'; + else if (devices[dn].has(aspect::host)) + ss << "host" << '\n'; + } catch (runtime_error re) { + // \todo handle errors + ss << "unknown\n"; + } + } + + std::cout << ss.str(); + ++i; + } } /*! diff --git a/backends/source/dppl_sycl_program_interface.cpp b/backends/source/dppl_sycl_program_interface.cpp new file mode 100644 index 0000000000..01d615a08b --- /dev/null +++ b/backends/source/dppl_sycl_program_interface.cpp @@ -0,0 +1,180 @@ +//===--- dppl_sycl_program_interface.cpp - DPPL-SYCL interface --*-- C++ -*-===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 the functions declared in +/// dppl_sycl_program_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dppl_sycl_program_interface.h" +#include "Support/CBindingWrapping.h" + +#include /* Sycl headers */ +#include /* OpenCL headers */ + +using namespace cl::sycl; + +namespace +{ +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPPLSyclContextRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(program, DPPLSyclProgramRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPPLSyclKernelRef) + +} /* end of anonymous namespace */ + +__dppl_give DPPLSyclProgramRef +DPPLProgram_CreateFromOCLSpirv (__dppl_keep const DPPLSyclContextRef CtxRef, + __dppl_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 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(); + + // Get a list of CL Devices from the Sycl devices + auto CLDevices = new cl_device_id[SyclDevices.size()]; + for (auto i = 0ul; i < SyclDevices.size(); ++i) + CLDevices[i] = SyclDevices[i].get(); + + // Build the OpenCL interoperability program + err = clBuildProgram(CLProgram, (cl_uint)(SyclDevices.size()), CLDevices, + nullptr, nullptr, nullptr); + // free the CLDevices array + delete[] CLDevices; + + if (err) { + // \todo: record the error string and any other information. + std::cerr << "OpenCL program could not be built. OpenCL Error " + << err << ".\n"; + return nullptr; + } + + // Create the Sycl program from OpenCL program + try { + auto SyclProgram = new program(*SyclCtx, CLProgram); + return wrap(SyclProgram); + } catch (invalid_object_error) { + // \todo record error + return nullptr; + } +} + +__dppl_give DPPLSyclProgramRef +DPPLProgram_CreateFromOCLSource (__dppl_keep const DPPLSyclContextRef Ctx, + __dppl_keep const char *Source, + __dppl_keep const char *CompileOpts) +{ + cl_int err; + std::string compileOpts; + context *SyclCtx = nullptr; + program *SyclProgram = nullptr; + + if(!Ctx) { + // \todo handle error + return nullptr; + } + + if(!Source) { + // \todo handle error message + return nullptr; + } + + SyclCtx = unwrap(Ctx); + SyclProgram = new program(*SyclCtx); + std::string source = Source; + + if(CompileOpts) { + compileOpts = CompileOpts; + } + + try{ + SyclProgram->build_with_source(source, compileOpts); + return wrap(SyclProgram); + } catch (compile_program_error) { + delete SyclProgram; + // \todo record error + return nullptr; + } catch (feature_not_supported) { + delete SyclProgram; + // \todo record error + return nullptr; + } +} + +__dppl_give DPPLSyclKernelRef +DPPLProgram_GetKernel (__dppl_keep DPPLSyclProgramRef PRef, + __dppl_keep const char *KernelName) +{ + if(!PRef) { + // \todo record error + return nullptr; + } + auto SyclProgram = unwrap(PRef); + if(!KernelName) { + // \todo record error + return nullptr; + } + std::string name = KernelName; + try { + auto SyclKernel = new kernel(SyclProgram->get_kernel(name)); + return wrap(SyclKernel); + } catch (invalid_object_error) { + // \todo record error + return nullptr; + } +} + +bool +DPPLProgram_HasKernel (__dppl_keep DPPLSyclProgramRef PRef, + __dppl_keep const char *KernelName) +{ + if(!PRef) { + // \todo handle error + return false; + } + + auto SyclProgram = unwrap(PRef); + try { + return SyclProgram->has_kernel(KernelName); + } catch (invalid_object_error) { + return false; + } +} + +void +DPPLProgram_Delete (__dppl_take DPPLSyclProgramRef PRef) +{ + delete unwrap(PRef); +} diff --git a/backends/source/dppl_sycl_queue_interface.cpp b/backends/source/dppl_sycl_queue_interface.cpp index c2f462badf..7ddaece07e 100644 --- a/backends/source/dppl_sycl_queue_interface.cpp +++ b/backends/source/dppl_sycl_queue_interface.cpp @@ -34,16 +34,95 @@ using namespace cl::sycl; namespace { // Create wrappers for C Binding types (see CBindingWrapping.h). - DEFINE_SIMPLE_CONVERSION_FUNCTIONS(queue, DPPLSyclQueueRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPPLSyclContextRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(device, DPPLSyclDeviceRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(event, DPPLSyclEventRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPPLSyclKernelRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(queue, DPPLSyclQueueRef) + +/*! + * @brief Set the kernel arg object + * + * @param cgh My Param doc + * @param Arg My Param doc + */ +bool set_kernel_arg (handler &cgh, size_t idx, __dppl_keep void *Arg, + DPPLKernelArgType ArgTy) +{ + bool arg_set = true; + + switch (ArgTy) + { + case DPPL_CHAR: + cgh.set_arg(idx, *(char*)Arg); + break; + case DPPL_SIGNED_CHAR: + cgh.set_arg(idx, *(signed char*)Arg); + break; + case DPPL_UNSIGNED_CHAR: + cgh.set_arg(idx, *(unsigned char*)Arg); + break; + case DPPL_SHORT: + cgh.set_arg(idx, *(short*)Arg); + break; + case DPPL_INT: + cgh.set_arg(idx, *(int*)Arg); + break; + case DPPL_UNSIGNED_INT: + cgh.set_arg(idx, *(unsigned int*)Arg); + break; + case DPPL_LONG: + cgh.set_arg(idx, *(long*)Arg); + break; + case DPPL_UNSIGNED_LONG: + cgh.set_arg(idx, *(unsigned long*)Arg); + break; + case DPPL_LONG_LONG: + cgh.set_arg(idx, *(long long*)Arg); + break; + case DPPL_UNSIGNED_LONG_LONG: + cgh.set_arg(idx, *(unsigned long long*)Arg); + break; + case DPPL_SIZE_T: + cgh.set_arg(idx, *(size_t*)Arg); + break; + case DPPL_FLOAT: + cgh.set_arg(idx, *(float*)Arg); + break; + case DPPL_DOUBLE: + cgh.set_arg(idx, *(double*)Arg); + break; + case DPPL_LONG_DOUBLE: + cgh.set_arg(idx, *(long double*)Arg); + break; + case DPPL_VOID_PTR: + cgh.set_arg(idx, Arg); + break; + default: + // \todo handle errors + arg_set = false; + std::cerr << "Kernel argument could not be created.\n"; + break; + } + return arg_set; +} } /* end of anonymous namespace */ +/*! + * Delete the passed in pointer after verifying it points to a sycl::queue. + */ +void DPPLQueue_Delete (__dppl_take DPPLSyclQueueRef QRef) +{ + delete unwrap(QRef); +} + __dppl_give DPPLSyclDeviceRef DPPLQueue_GetDevice (__dppl_keep const DPPLSyclQueueRef QRef) { auto Q = unwrap(QRef); auto Device = new device(Q->get_device()); - return reinterpret_cast(Device); + return wrap(Device); } __dppl_give DPPLSyclContextRef @@ -51,15 +130,134 @@ DPPLQueue_GetContext (__dppl_keep const DPPLSyclQueueRef QRef) { auto Q = unwrap(QRef); auto Context = new context(Q->get_context()); - return reinterpret_cast(Context); + return wrap(Context); } -/*! - * Delete the passed in pointer after verifying it points to a sycl::queue. - */ -void DPPLQueue_Delete (__dppl_take DPPLSyclQueueRef QRef) +__dppl_give DPPLSyclEventRef +DPPLQueue_SubmitRange (__dppl_keep const DPPLSyclKernelRef KRef, + __dppl_keep const DPPLSyclQueueRef QRef, + __dppl_keep void **Args, + __dppl_keep const DPPLKernelArgType *ArgTypes, + size_t NArgs, + __dppl_keep const size_t Range[3], + size_t NDims, + __dppl_keep const DPPLSyclEventRef *DepEvents, + size_t NDepEvents) { - delete unwrap(QRef); + auto Kernel = unwrap(KRef); + auto Queue = unwrap(QRef); + event e; + + try { + e = Queue->submit([&](handler& cgh) { + // Depend on any event that was specified by the caller. + if(NDepEvents) + for(auto i = 0ul; i < NDepEvents; ++i) + cgh.depends_on(*unwrap(DepEvents[i])); + + for (auto i = 0ul; i < NArgs; ++i) { + // \todo add support for Sycl buffers + // \todo handle errors properly + if(!set_kernel_arg(cgh, i, Args[i], ArgTypes[i])) + exit(1); + } + switch(NDims) + { + case 1: + cgh.parallel_for(range<1>{Range[0]}, *Kernel); + break; + case 2: + cgh.parallel_for(range<2>{Range[0], Range[1]}, *Kernel); + break; + case 3: + cgh.parallel_for(range<3>{Range[0], Range[1], Range[2]}, + *Kernel); + break; + default: + // \todo handle the error + throw std::runtime_error("Range cannot be greater than three " + "dimensions."); + } + }); + } catch (runtime_error re) { + // \todo fix error handling + std::cerr << re.what() << '\n'; + return nullptr; + } catch (std::runtime_error sre) { + std::cerr << sre.what() << '\n'; + return nullptr; + } + + return wrap(new event(e)); +} + +DPPLSyclEventRef +DPPLQueue_SubmitNDRange(__dppl_keep const DPPLSyclKernelRef KRef, + __dppl_keep const DPPLSyclQueueRef QRef, + __dppl_keep void **Args, + __dppl_keep const DPPLKernelArgType *ArgTypes, + size_t NArgs, + __dppl_keep const size_t gRange[3], + __dppl_keep const size_t lRange[3], + size_t NDims, + __dppl_keep const DPPLSyclEventRef *DepEvents, + size_t NDepEvents) +{ + auto Kernel = unwrap(KRef); + auto Queue = unwrap(QRef); + event e; + + try { + e = Queue->submit([&](handler& cgh) { + // Depend on any event that was specified by the caller. + if(NDepEvents) + for(auto i = 0ul; i < NDepEvents; ++i) + cgh.depends_on(*unwrap(DepEvents[i])); + + for (auto i = 0ul; i < NArgs; ++i) { + // \todo add support for Sycl buffers + // \todo handle errors properly + if(!set_kernel_arg(cgh, i, Args[i], ArgTypes[i])) + exit(1); + } + switch(NDims) + { + case 1: + cgh.parallel_for(nd_range<1>{{gRange[0]},{lRange[0]}}, *Kernel); + break; + case 2: + cgh.parallel_for(nd_range<2>{{gRange[0], gRange[1]}, + {lRange[0], lRange[1]}}, *Kernel); + break; + case 3: + cgh.parallel_for(nd_range<3>{{gRange[0], gRange[1], gRange[2]}, + {lRange[0], lRange[1], lRange[3]}}, + *Kernel); + break; + default: + // \todo handle the error + throw std::runtime_error("Range cannot be greater than three " + "dimensions."); + } + }); + } catch (runtime_error re) { + // \todo fix error handling + std::cerr << re.what() << '\n'; + return nullptr; + } catch (std::runtime_error sre) { + std::cerr << sre.what() << '\n'; + return nullptr; + } + + return wrap(new event(e)); +} + +void +DPPLQueue_Wait (__dppl_keep DPPLSyclQueueRef QRef) +{ + // \todo what happens if the QRef is null or a pointer to a valid sycl queue + auto SyclQueue = unwrap(QRef); + SyclQueue->wait(); } void DPPLQueue_Memcpy (__dppl_take const DPPLSyclQueueRef QRef, diff --git a/backends/source/dppl_sycl_queue_manager.cpp b/backends/source/dppl_sycl_queue_manager.cpp index 3fa6033016..3981c528dc 100644 --- a/backends/source/dppl_sycl_queue_manager.cpp +++ b/backends/source/dppl_sycl_queue_manager.cpp @@ -39,8 +39,8 @@ using namespace cl::sycl; namespace { - // Create wrappers for C Binding types (see CBindingWrapping.h). - DEFINE_SIMPLE_CONVERSION_FUNCTIONS(queue, DPPLSyclQueueRef) +// Create wrappers for C Binding types (see CBindingWrapping.h). +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(queue, DPPLSyclQueueRef) void error_reporter (const std::string & msg) { @@ -99,6 +99,7 @@ class QMgrHelper static cl::sycl::vector_class* init_queues (info::device_type device_ty) { + auto queues = new std::vector(); for(auto d : device::get_devices(device_ty)) queues->emplace_back(d); diff --git a/backends/source/dppl_sycl_usm_interface.cpp b/backends/source/dppl_sycl_usm_interface.cpp index 4e1444634b..54ff95efc6 100644 --- a/backends/source/dppl_sycl_usm_interface.cpp +++ b/backends/source/dppl_sycl_usm_interface.cpp @@ -34,9 +34,9 @@ using namespace cl::sycl; namespace { // Create wrappers for C Binding types (see CBindingWrapping.h). - DEFINE_SIMPLE_CONVERSION_FUNCTIONS(queue, DPPLSyclQueueRef) - DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPPLSyclContextRef) - DEFINE_SIMPLE_CONVERSION_FUNCTIONS(void, DPPLSyclUSMRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(queue, DPPLSyclQueueRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPPLSyclContextRef) +DEFINE_SIMPLE_CONVERSION_FUNCTIONS(void, DPPLSyclUSMRef) } /* end of anonymous namespace */ diff --git a/backends/tests/CMakeLists.txt b/backends/tests/CMakeLists.txt index c3add176ba..08e7f9f9ff 100644 --- a/backends/tests/CMakeLists.txt +++ b/backends/tests/CMakeLists.txt @@ -23,10 +23,21 @@ else() link_directories(${GTEST_LIB_DIR}) set(PYDPPL_BACKEND_TEST_CASES - test_sycl_queue_manager + test_sycl_kernel_interface test_sycl_platform_interface + test_sycl_program_interface + test_sycl_queue_interface + test_sycl_queue_manager ) + # Copy the spir-v input files to test build directory + set(spirv-test-files + multi_kernel.spv + ) + foreach(tf ${spirv-test-files}) + file(COPY ${tf} DESTINATION ${CMAKE_CURRENT_BINARY_DIR}) + endforeach() + foreach(TEST_CASE ${PYDPPL_BACKEND_TEST_CASES}) add_executable(${TEST_CASE} EXCLUDE_FROM_ALL ${TEST_CASE}.cpp) target_link_libraries( diff --git a/backends/tests/multi_kernel.spv b/backends/tests/multi_kernel.spv new file mode 100644 index 0000000000..1fa772d8df Binary files /dev/null and b/backends/tests/multi_kernel.spv differ diff --git a/backends/tests/test_sycl_kernel_interface.cpp b/backends/tests/test_sycl_kernel_interface.cpp new file mode 100644 index 0000000000..5a360dab62 --- /dev/null +++ b/backends/tests/test_sycl_kernel_interface.cpp @@ -0,0 +1,106 @@ +//===---- test_sycl_program_interface.cpp - DPPL-SYCL interface -*- C++ -*-===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 has unit test cases for functions defined in +/// dppl_sycl_kernel_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dppl_sycl_context_interface.h" +#include "dppl_sycl_kernel_interface.h" +#include "dppl_sycl_program_interface.h" +#include "dppl_sycl_queue_interface.h" +#include "dppl_sycl_queue_manager.h" +#include "dppl_utils.h" + +#include +#include +#include + +using namespace cl::sycl; + + +struct TestDPPLSyclKernelInterface : public ::testing::Test +{ + const char *CLProgramStr = R"CLC( + kernel void add(global int* a, global int* b, global int* c) { + size_t index = get_global_id(0); + c[index] = a[index] + b[index]; + } + + kernel void axpy(global int* a, global int* b, global int* c, int d) { + size_t index = get_global_id(0); + c[index] = a[index] + d*b[index]; + } + )CLC"; + const char *CompileOpts ="-cl-fast-relaxed-math"; + + DPPLSyclContextRef CtxRef = nullptr; + DPPLSyclQueueRef QueueRef = nullptr; + DPPLSyclProgramRef PRef = nullptr; + DPPLSyclKernelRef AddKernel = nullptr; + DPPLSyclKernelRef AxpyKernel = nullptr; + + TestDPPLSyclKernelInterface () + { + QueueRef = DPPLQueueMgr_GetQueue(DPPL_GPU, 0); + CtxRef = DPPLQueue_GetContext(QueueRef); + PRef = DPPLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, + CompileOpts); + AddKernel = DPPLProgram_GetKernel(PRef, "add"); + AxpyKernel = DPPLProgram_GetKernel(PRef, "axpy"); + } + + ~TestDPPLSyclKernelInterface () + { + DPPLQueue_Delete(QueueRef); + DPPLContext_Delete(CtxRef); + DPPLProgram_Delete(PRef); + DPPLKernel_Delete(AddKernel); + DPPLKernel_Delete(AxpyKernel); + } +}; + +TEST_F (TestDPPLSyclKernelInterface, CheckGetFunctionName) +{ + + auto fnName1 = DPPLKernel_GetFunctionName(AddKernel); + auto fnName2 = DPPLKernel_GetFunctionName(AxpyKernel); + ASSERT_STREQ("add", fnName1); + ASSERT_STREQ("axpy", fnName2); + DPPLCString_Delete(fnName1); + DPPLCString_Delete(fnName2); +} + +TEST_F (TestDPPLSyclKernelInterface, CheckGetNumArgs) +{ + + ASSERT_EQ(DPPLKernel_GetNumArgs(AddKernel), 3); + ASSERT_EQ(DPPLKernel_GetNumArgs(AxpyKernel), 4); +} + +int +main (int argc, char** argv) +{ + ::testing::InitGoogleTest(&argc, argv); + int ret = RUN_ALL_TESTS(); + return ret; +} diff --git a/backends/tests/test_sycl_program_interface.cpp b/backends/tests/test_sycl_program_interface.cpp new file mode 100644 index 0000000000..3b2fff423b --- /dev/null +++ b/backends/tests/test_sycl_program_interface.cpp @@ -0,0 +1,212 @@ +//===---- test_sycl_program_interface.cpp - DPPL-SYCL interface -*- C++ -*-===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 has unit test cases for functions defined in +/// dppl_sycl_program_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dppl_sycl_context_interface.h" +#include "dppl_sycl_kernel_interface.h" +#include "dppl_sycl_program_interface.h" +#include "dppl_sycl_queue_interface.h" +#include "dppl_sycl_queue_manager.h" + +#include +#include +#include +#include +#include + +using namespace cl::sycl; + +namespace +{ + const size_t SIZE = 1024; + + void add_kernel_checker (queue *syclQueue, DPPLSyclKernelRef AddKernel) + { + range<1> a_size{SIZE}; + std::array a, b, c; + + for (int i = 0; i a_device(a.data(), a_size); + buffer b_device(b.data(), a_size); + buffer c_device(c.data(), a_size); + buffer buffs[3] = {a_device, b_device, c_device}; + syclQueue->submit([&](handler& cgh) { + for (auto buff : buffs) { + auto arg = buff.get_access(cgh); + cgh.set_args(arg); + } + auto syclKernel = reinterpret_cast(AddKernel); + cgh.parallel_for(range<1>{SIZE}, *syclKernel); + }); + } + + // Validate the data + for(auto i = 0ul; i < SIZE; ++i) { + EXPECT_EQ(c[i], i + i); + } + } + + void axpy_kernel_checker (queue *syclQueue, DPPLSyclKernelRef AxpyKernel) + { + range<1> a_size{SIZE}; + std::array a, b, c; + + for (int i = 0; i a_device(a.data(), a_size); + buffer b_device(b.data(), a_size); + buffer c_device(c.data(), a_size); + buffer buffs[3] = {a_device, b_device, c_device}; + syclQueue->submit([&](handler& cgh) { + for (auto i = 0ul; i < 3; ++i) { + auto arg = buffs[i].get_access(cgh); + cgh.set_arg(i, arg); + } + cgh.set_arg(3, d); + auto syclKernel = reinterpret_cast(AxpyKernel); + cgh.parallel_for(range<1>{SIZE}, *syclKernel); + }); + } + + // Validate the data + for(auto i = 0ul; i < SIZE; ++i) { + EXPECT_EQ(c[i], i + d*i); + } + } +} + +struct TestDPPLSyclProgramInterface : public ::testing::Test +{ + const char *CLProgramStr = R"CLC( + kernel void add(global int* a, global int* b, global int* c) { + size_t index = get_global_id(0); + c[index] = a[index] + b[index]; + } + + kernel void axpy(global int* a, global int* b, global int* c, int d) { + size_t index = get_global_id(0); + c[index] = a[index] + d*b[index]; + } + )CLC"; + const char *CompileOpts ="-cl-fast-relaxed-math"; + + DPPLSyclContextRef CtxRef = nullptr; + DPPLSyclQueueRef QueueRef = nullptr; + DPPLSyclProgramRef PRef = nullptr; + DPPLSyclProgramRef PRef2 = nullptr; + + TestDPPLSyclProgramInterface () + { + QueueRef = DPPLQueueMgr_GetQueue(DPPL_GPU, 0); + CtxRef = DPPLQueue_GetContext(QueueRef); + PRef = DPPLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, + CompileOpts); + + // Create a program from a SPIR-V file + std::ifstream file{"./multi_kernel.spv", + std::ios::binary | std::ios::ate}; + auto fileSize = std::filesystem::file_size("./multi_kernel.spv"); + file.seekg(0, std::ios::beg); + std::vector buffer(fileSize); + file.read(buffer.data(), fileSize); + PRef2 = DPPLProgram_CreateFromOCLSpirv(CtxRef, buffer.data(), + fileSize); + } + + ~TestDPPLSyclProgramInterface () + { + DPPLQueue_Delete(QueueRef); + DPPLContext_Delete(CtxRef); + DPPLProgram_Delete(PRef); + DPPLProgram_Delete(PRef2); + } +}; + +TEST_F (TestDPPLSyclProgramInterface, CheckCreateFromOCLSource) +{ + ASSERT_TRUE(PRef != nullptr); +} + +TEST_F (TestDPPLSyclProgramInterface, CheckCreateFromOCLSpirv) +{ + ASSERT_TRUE(PRef2 != nullptr); +} + +TEST_F (TestDPPLSyclProgramInterface, CheckHasKernelOCLSource) +{ + ASSERT_TRUE(DPPLProgram_HasKernel(PRef, "add")); + ASSERT_TRUE(DPPLProgram_HasKernel(PRef, "axpy")); +} + +TEST_F (TestDPPLSyclProgramInterface, CheckHasKernelSpirvSource) +{ + ASSERT_TRUE(DPPLProgram_HasKernel(PRef, "add")); + ASSERT_TRUE(DPPLProgram_HasKernel(PRef, "axpy")); +} + +TEST_F (TestDPPLSyclProgramInterface, CheckGetKernelOCLSource) +{ + auto AddKernel = DPPLProgram_GetKernel(PRef, "add"); + auto AxpyKernel = DPPLProgram_GetKernel(PRef, "axpy"); + auto syclQueue = reinterpret_cast(QueueRef); + + add_kernel_checker(syclQueue, AddKernel); + axpy_kernel_checker(syclQueue, AxpyKernel); + + DPPLKernel_Delete(AddKernel); + DPPLKernel_Delete(AxpyKernel); +} + +TEST_F (TestDPPLSyclProgramInterface, CheckGetKernelOCLSpirv) +{ + auto AddKernel = DPPLProgram_GetKernel(PRef2, "add"); + auto AxpyKernel = DPPLProgram_GetKernel(PRef2, "axpy"); + auto syclQueue = reinterpret_cast(QueueRef); + + add_kernel_checker(syclQueue, AddKernel); + axpy_kernel_checker(syclQueue, AxpyKernel); + + DPPLKernel_Delete(AddKernel); + DPPLKernel_Delete(AxpyKernel); +} + +int +main (int argc, char** argv) +{ + ::testing::InitGoogleTest(&argc, argv); + int ret = RUN_ALL_TESTS(); + return ret; +} diff --git a/backends/tests/test_sycl_queue_interface.cpp b/backends/tests/test_sycl_queue_interface.cpp new file mode 100644 index 0000000000..d284799540 --- /dev/null +++ b/backends/tests/test_sycl_queue_interface.cpp @@ -0,0 +1,192 @@ +//===---- test_sycl_queue_interface.cpp - DPPL-SYCL interface -*- C++ --*--===// +// +// Python Data Parallel Processing Library (PyDPPL) +// +// 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 has unit test cases for functions defined in +/// dppl_sycl_queue_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dppl_sycl_context_interface.h" +#include "dppl_sycl_event_interface.h" +#include "dppl_sycl_kernel_interface.h" +#include "dppl_sycl_program_interface.h" +#include "dppl_sycl_queue_interface.h" +#include "dppl_sycl_queue_manager.h" +#include "dppl_sycl_usm_interface.h" + +#include "Support/CBindingWrapping.h" + +#include + +namespace +{ + constexpr size_t SIZE = 1024; + + DEFINE_SIMPLE_CONVERSION_FUNCTIONS(void, DPPLSyclUSMRef); + + void add_kernel_checker (const float *a, const float *b, const float *c) + { + // Validate the data + for(auto i = 0ul; i < SIZE; ++i) { + EXPECT_EQ(c[i], a[i] + b[i]); + } + } + + void axpy_kernel_checker (const float *a, const float *b, const float *c, + float d) + { + for(auto i = 0ul; i < SIZE; ++i) { + EXPECT_EQ(c[i], a[i] + d*b[i]); + } + } +} + +struct TestDPPLSyclQueueInterface : public ::testing::Test +{ + const char *CLProgramStr = R"CLC( + kernel void init_arr (global float *a) { + size_t index = get_global_id(0); + a[index] = (float)index; + } + + kernel void add (global float* a, global float* b, global float* c) { + size_t index = get_global_id(0); + c[index] = a[index] + b[index]; + } + + kernel void axpy (global float* a, global float* b, + global float* c, float d) { + size_t index = get_global_id(0); + c[index] = a[index] + d*b[index]; + } + )CLC"; + const char *CompileOpts ="-cl-fast-relaxed-math"; + + DPPLSyclContextRef CtxRef = nullptr; + DPPLSyclQueueRef Queue = nullptr; + DPPLSyclProgramRef PRef = nullptr; + DPPLSyclProgramRef PRef2 = nullptr; + TestDPPLSyclQueueInterface () + { + Queue = DPPLQueueMgr_GetQueue(DPPL_GPU, 0); + CtxRef = DPPLQueue_GetContext(Queue); + PRef = DPPLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, + CompileOpts); + } + + ~TestDPPLSyclQueueInterface () + { + DPPLQueue_Delete(Queue); + DPPLContext_Delete(CtxRef); + DPPLProgram_Delete(PRef); + } +}; + +TEST_F (TestDPPLSyclQueueInterface, CheckSubmit) +{ + ASSERT_TRUE(PRef != nullptr); + ASSERT_TRUE(DPPLProgram_HasKernel(PRef, "init_arr")); + ASSERT_TRUE(DPPLProgram_HasKernel(PRef, "add")); + ASSERT_TRUE(DPPLProgram_HasKernel(PRef, "axpy")); + + auto InitKernel = DPPLProgram_GetKernel(PRef, "init_arr"); + auto AddKernel = DPPLProgram_GetKernel(PRef, "add"); + auto AxpyKernel = DPPLProgram_GetKernel(PRef, "axpy"); + + // Create the input args + auto a = DPPLmalloc_shared(SIZE, Queue); + ASSERT_TRUE(a != nullptr); + auto b = DPPLmalloc_shared(SIZE, Queue); + ASSERT_TRUE(b != nullptr); + auto c = DPPLmalloc_shared(SIZE, Queue); + ASSERT_TRUE(c != nullptr); + + // Initialize a,b + DPPLKernelArgType argTypes[] = {DPPL_VOID_PTR}; + size_t Range[] = {SIZE}; + void *arg1[1] = { unwrap(a) }; + void *arg2[1] = { unwrap(b) }; + + auto E1 = DPPLQueue_SubmitRange(InitKernel, Queue, arg1, argTypes, 1, + Range, 1, nullptr, 0); + auto E2 = DPPLQueue_SubmitRange(InitKernel, Queue, arg2, argTypes, 1, + Range, 1, nullptr, 0); + ASSERT_TRUE(E1 != nullptr); + ASSERT_TRUE(E2 != nullptr); + + DPPLQueue_Wait(Queue); + + // Submit the add kernel + void *args[3] = { unwrap(a), unwrap(b), unwrap(c) }; + DPPLKernelArgType addKernelArgTypes[] = { + DPPL_VOID_PTR, + DPPL_VOID_PTR, + DPPL_VOID_PTR + }; + + auto E3 = DPPLQueue_SubmitRange(AddKernel, Queue, args, + addKernelArgTypes, 3, Range, 1, nullptr, 0); + ASSERT_TRUE(E3 != nullptr); + DPPLQueue_Wait(Queue); + + // Verify the result of "add" + add_kernel_checker((float*)a, (float*)b, (float*)c); + + // Create kernel args for axpy + float d = 10.0; + void *args2[4] = { unwrap(a), unwrap(b), unwrap(c) , (void*)&d }; + DPPLKernelArgType addKernelArgTypes2[] = { + DPPL_VOID_PTR, + DPPL_VOID_PTR, + DPPL_VOID_PTR, + DPPL_FLOAT + }; + auto E4 = DPPLQueue_SubmitRange(AxpyKernel, Queue, args2, + addKernelArgTypes2, 4, Range, 1, + nullptr, 0); + ASSERT_TRUE(E4 != nullptr); + DPPLQueue_Wait(Queue); + + // Verify the result of "axpy" + axpy_kernel_checker((float*)a, (float*)b, (float*)c, d); + + // clean ups + DPPLEvent_Delete(E1); + DPPLEvent_Delete(E2); + DPPLEvent_Delete(E3); + DPPLEvent_Delete(E4); + + DPPLKernel_Delete(AddKernel); + DPPLKernel_Delete(AxpyKernel); + DPPLKernel_Delete(InitKernel); + + DPPLfree_with_queue((DPPLSyclUSMRef)a, Queue); + DPPLfree_with_queue((DPPLSyclUSMRef)b, Queue); + DPPLfree_with_queue((DPPLSyclUSMRef)c, Queue); +} + +int +main (int argc, char** argv) +{ + ::testing::InitGoogleTest(&argc, argv); + int ret = RUN_ALL_TESTS(); + return ret; +} diff --git a/backends/tests/test_sycl_queue_manager.cpp b/backends/tests/test_sycl_queue_manager.cpp index 05b04aabf6..38faadc217 100644 --- a/backends/tests/test_sycl_queue_manager.cpp +++ b/backends/tests/test_sycl_queue_manager.cpp @@ -55,11 +55,11 @@ namespace } } -struct TestDPPLSyclQueuemanager : public ::testing::Test +struct TestDPPLSyclQueueManager : public ::testing::Test { }; -TEST_F (TestDPPLSyclQueuemanager, CheckDPPLGetCurrentQueue) +TEST_F (TestDPPLSyclQueueManager, CheckDPPLGetCurrentQueue) { DPPLSyclQueueRef q; ASSERT_NO_THROW(q = DPPLQueueMgr_GetCurrentQueue()); @@ -67,7 +67,7 @@ TEST_F (TestDPPLSyclQueuemanager, CheckDPPLGetCurrentQueue) } -TEST_F (TestDPPLSyclQueuemanager, CheckDPPLGetQueue) +TEST_F (TestDPPLSyclQueueManager, CheckDPPLGetQueue) { auto numCpuQueues = DPPLQueueMgr_GetNumCPUQueues(); auto numGpuQueues = DPPLQueueMgr_GetNumGPUQueues(); @@ -94,7 +94,7 @@ TEST_F (TestDPPLSyclQueuemanager, CheckDPPLGetQueue) } -TEST_F (TestDPPLSyclQueuemanager, CheckGetNumActivatedQueues) +TEST_F (TestDPPLSyclQueueManager, CheckGetNumActivatedQueues) { size_t num0, num1, num2, num4; @@ -124,7 +124,7 @@ TEST_F (TestDPPLSyclQueuemanager, CheckGetNumActivatedQueues) } -TEST_F (TestDPPLSyclQueuemanager, CheckDPPLDumpDeviceInfo) +TEST_F (TestDPPLSyclQueueManager, CheckDPPLDumpDeviceInfo) { auto q = DPPLQueueMgr_GetCurrentQueue(); EXPECT_NO_FATAL_FAILURE(DPPLDevice_DumpInfo(DPPLQueue_GetDevice(q))); diff --git a/dpctl/.gitignore b/dpctl/.gitignore index 1edc29f09d..957638d06f 100644 --- a/dpctl/.gitignore +++ b/dpctl/.gitignore @@ -1,3 +1,4 @@ *.so *.cpp -*.c \ No newline at end of file +*.c +include/* diff --git a/dpctl/_sycl_core.pxd b/dpctl/_sycl_core.pxd index b10afa787e..c8860b3c6b 100644 --- a/dpctl/_sycl_core.pxd +++ b/dpctl/_sycl_core.pxd @@ -33,7 +33,7 @@ from .backend cimport * cdef class SyclContext: ''' Wrapper class for a Sycl Context ''' - cdef DPPLSyclContextRef ctxt_ptr + cdef DPPLSyclContextRef _ctxt_ref @staticmethod cdef SyclContext _create (DPPLSyclContextRef ctxt) @@ -43,24 +43,76 @@ cdef class SyclContext: cdef class SyclDevice: ''' Wrapper class for a Sycl Device ''' - cdef DPPLSyclDeviceRef device_ptr - cdef const char *vendor_name - cdef const char *device_name - cdef const char *driver_version + cdef DPPLSyclDeviceRef _device_ref + cdef const char *_vendor_name + cdef const char *_device_name + cdef const char *_driver_version @staticmethod cdef SyclDevice _create (DPPLSyclDeviceRef dref) - cdef DPPLSyclDeviceRef get_device_ptr (self) + cdef DPPLSyclDeviceRef get_device_ref (self) + + +cdef class SyclEvent: + ''' Wrapper class for a Sycl Event + ''' + cdef DPPLSyclEventRef _event_ref + cdef list _args + + @staticmethod + cdef SyclEvent _create (DPPLSyclEventRef e, list args) + cdef DPPLSyclEventRef get_event_ref (self) + cpdef void wait (self) + + +cdef class SyclKernel: + ''' Wraps a sycl::kernel object created from an OpenCL interoperability + kernel. + ''' + cdef DPPLSyclKernelRef _kernel_ref + cdef const char *_function_name + cdef DPPLSyclKernelRef get_kernel_ref (self) + + @staticmethod + cdef SyclKernel _create (DPPLSyclKernelRef kref) + + +cdef class SyclProgram: + ''' Wraps a sycl::program object created from an OpenCL interoperability + program. + + SyclProgram exposes the C API from dppl_sycl_program_interface.h. A + SyclProgram can be created from either a source string or a SPIR-V + binary file. + ''' + cdef DPPLSyclProgramRef _program_ref + + @staticmethod + cdef SyclProgram _create (DPPLSyclProgramRef pref) + cdef DPPLSyclProgramRef get_program_ref (self) + cpdef SyclKernel get_sycl_kernel(self, str kernel_name) cdef class SyclQueue: ''' Wrapper class for a Sycl queue. ''' - cdef DPPLSyclQueueRef queue_ptr + cdef DPPLSyclQueueRef _queue_ref + cdef SyclContext _context + cdef SyclDevice _device + + cdef _raise_queue_submit_error (self, fname, errcode) + cdef _raise_invalid_range_error (self, fname, ndims, errcode) + cdef int _populate_args (self, list args, void **kargs, + DPPLKernelArgType *kargty) + cdef int _populate_range (self, size_t Range[3], list gS, size_t nGS) @staticmethod - cdef SyclQueue _create (DPPLSyclQueueRef qref) + cdef SyclQueue _create (DPPLSyclQueueRef qref) cpdef SyclContext get_sycl_context (self) cpdef SyclDevice get_sycl_device (self) + cdef DPPLSyclQueueRef get_queue_ref (self) + cpdef SyclEvent submit (self, SyclKernel kernel, list args, list gS, + list lS=*, list dEvents=*) + cpdef void wait (self) cdef DPPLSyclQueueRef get_queue_ref (self) cpdef memcpy (self, dest, src, int count) diff --git a/dpctl/backend.pxd b/dpctl/backend.pxd index cab1b7430d..3f587b77fe 100644 --- a/dpctl/backend.pxd +++ b/dpctl/backend.pxd @@ -36,14 +36,20 @@ cdef extern from "dppl_utils.h": cdef extern from "dppl_sycl_types.h": cdef struct DPPLOpaqueSyclContext - cdef struct DPPLOpaqueSyclQueue cdef struct DPPLOpaqueSyclDevice + cdef struct DPPLOpaqueSyclEvent + cdef struct DPPLOpaqueSyclKernel + cdef struct DPPLOpaqueSyclProgram + cdef struct DPPLOpaqueSyclQueue cdef struct DPPLOpaqueSyclUSM ctypedef DPPLOpaqueSyclContext* DPPLSyclContextRef - ctypedef DPPLOpaqueSyclQueue* DPPLSyclQueueRef - ctypedef DPPLOpaqueSyclDevice* DPPLSyclDeviceRef - ctypedef DPPLOpaqueSyclUSM* DPPLSyclUSMRef + ctypedef DPPLOpaqueSyclDevice* DPPLSyclDeviceRef + ctypedef DPPLOpaqueSyclEvent* DPPLSyclEventRef + ctypedef DPPLOpaqueSyclKernel* DPPLSyclKernelRef + ctypedef DPPLOpaqueSyclProgram* DPPLSyclProgramRef + ctypedef DPPLOpaqueSyclQueue* DPPLSyclQueueRef + ctypedef DPPLOpaqueSyclUSM* DPPLSyclUSMRef cdef extern from "dppl_sycl_context_interface.h": @@ -51,6 +57,9 @@ cdef extern from "dppl_sycl_context_interface.h": cdef extern from "dppl_sycl_device_interface.h": + cdef enum _device_type 'DPPLSyclDeviceType': + _GPU 'DPPL_GPU' + _CPU 'DPPL_CPU' cdef void DPPLDevice_DumpInfo (const DPPLSyclDeviceRef DRef) except + cdef void DPPLDevice_Delete (DPPLSyclDeviceRef DRef) except + cdef void DPPLDevice_DumpInfo (const DPPLSyclDeviceRef DRef) except + @@ -67,27 +76,86 @@ cdef extern from "dppl_sycl_device_interface.h": except + +cdef extern from "dppl_sycl_event_interface.h": + cdef void DPPLEvent_Wait (DPPLSyclEventRef ERef) + cdef void DPPLEvent_Delete (DPPLSyclEventRef ERef) + + +cdef extern from "dppl_sycl_kernel_interface.h": + cdef const char* DPPLKernel_GetFunctionName (const DPPLSyclKernelRef KRef) + cdef size_t DPPLKernel_GetNumArgs (const DPPLSyclKernelRef KRef) + cdef void DPPLKernel_Delete (DPPLSyclKernelRef KRef) + + cdef extern from "dppl_sycl_platform_interface.h": cdef size_t DPPLPlatform_GetNumPlatforms () cdef void DPPLPlatform_DumpInfo () +cdef extern from "dppl_sycl_program_interface.h": + cdef DPPLSyclProgramRef DPPLProgram_CreateFromOCLSpirv ( + const DPPLSyclContextRef Ctx, + const void *IL, + size_t Length) + cdef DPPLSyclProgramRef DPPLProgram_CreateFromOCLSource ( + const DPPLSyclContextRef Ctx, + const char* Source, + const char* CompileOpts) + cdef DPPLSyclKernelRef DPPLProgram_GetKernel (DPPLSyclProgramRef PRef, + const char *KernelName) + cdef bool DPPLProgram_HasKernel (DPPLSyclProgramRef PRef, + const char *KernelName) + cdef void DPPLProgram_Delete (DPPLSyclProgramRef PRef) + + cdef extern from "dppl_sycl_queue_interface.h": - cdef void DPPLQueue_Delete (DPPLSyclQueueRef QRef) except + - cdef DPPLSyclContextRef DPPLQueue_GetContext (const DPPLSyclQueueRef Q) \ - except+ - cdef DPPLSyclDeviceRef DPPLQueue_GetDevice (const DPPLSyclQueueRef Q) \ - except + + cdef enum _arg_data_type 'DPPLKernelArgType': + _CHAR 'DPPL_CHAR', + _SIGNED_CHAR 'DPPL_SIGNED_CHAR', + _UNSIGNED_CHAR 'DPPL_UNSIGNED_CHAR', + _SHORT 'DPPL_SHORT', + _INT 'DPPL_INT', + _UNSIGNED_INT 'DPPL_INT', + _LONG 'DPPL_LONG', + _UNSIGNED_LONG 'DPPL_UNSIGNED_LONG', + _LONG_LONG 'DPPL_LONG_LONG', + _UNSIGNED_LONG_LONG 'DPPL_UNSIGNED_LONG_LONG', + _SIZE_T 'DPPL_SIZE_T', + _FLOAT 'DPPL_FLOAT', + _DOUBLE 'DPPL_DOUBLE', + _LONG_DOUBLE 'DPPL_DOUBLE', + _VOID_PTR 'DPPL_VOID_PTR' + ctypedef _arg_data_type DPPLKernelArgType + cdef void DPPLQueue_Delete (DPPLSyclQueueRef QRef) + cdef DPPLSyclContextRef DPPLQueue_GetContext (const DPPLSyclQueueRef Q) + cdef DPPLSyclDeviceRef DPPLQueue_GetDevice (const DPPLSyclQueueRef Q) + cdef DPPLSyclEventRef DPPLQueue_SubmitRange ( + const DPPLSyclKernelRef Ref, + const DPPLSyclQueueRef QRef, + void **Args, + const DPPLKernelArgType *ArgTypes, + size_t NArgs, + const size_t Range[3], + size_t NDims, + const DPPLSyclEventRef *DepEvents, + size_t NDepEvents) + cdef DPPLSyclEventRef DPPLQueue_SubmitNDRange( + const DPPLSyclKernelRef Ref, + const DPPLSyclQueueRef QRef, + void **Args, + const DPPLKernelArgType *ArgTypes, + size_t NArgs, + const size_t gRange[3], + const size_t lRange[3], + size_t NDims, + const DPPLSyclEventRef *DepEvents, + size_t NDepEvents) + cdef void DPPLQueue_Wait (const DPPLSyclQueueRef QRef) cdef void DPPLQueue_Memcpy (const DPPLSyclQueueRef Q, - void *Dest, const void *Src, size_t Count) \ - except + + void *Dest, const void *Src, size_t Count) cdef extern from "dppl_sycl_queue_manager.h": - cdef enum _device_type 'DPPLSyclDeviceType': - _GPU 'DPPL_GPU' - _CPU 'DPPL_CPU' - cdef DPPLSyclQueueRef DPPLQueueMgr_GetCurrentQueue () except + cdef size_t DPPLQueueMgr_GetNumCPUQueues () except + cdef size_t DPPLQueueMgr_GetNumGPUQueues () except + diff --git a/dpctl/sycl_core.pyx b/dpctl/sycl_core.pyx index b1a5df9b8c..347a25d29d 100644 --- a/dpctl/sycl_core.pyx +++ b/dpctl/sycl_core.pyx @@ -31,6 +31,7 @@ from enum import Enum, auto import logging from .backend cimport * from ._memory cimport Memory +from libc.stdlib cimport malloc, free _logger = logging.getLogger(__name__) @@ -41,26 +42,44 @@ class device_type(Enum): cpu = auto() -cdef class UnsupportedDeviceTypeError(Exception): +cdef class UnsupportedDeviceTypeError (Exception): '''This exception is raised when a device type other than CPU or GPU is encountered. ''' 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. + ''' + pass + +cdef class SyclKernelInvalidRangeError (Exception): + '''This exception is raised when a range that has more than three + dimensions or less than one dimension. + ''' + pass + cdef class SyclContext: @staticmethod cdef SyclContext _create (DPPLSyclContextRef ctxt): cdef SyclContext ret = SyclContext.__new__(SyclContext) - ret.ctxt_ptr = ctxt + ret._ctxt_ref = ctxt return ret def __dealloc__ (self): - DPPLContext_Delete(self.ctxt_ptr) + DPPLContext_Delete(self._ctxt_ref) cdef DPPLSyclContextRef get_context_ref (self): - return self.ctxt_ptr + return self._ctxt_ref cdef class SyclDevice: @@ -70,32 +89,32 @@ cdef class SyclDevice: @staticmethod cdef SyclDevice _create (DPPLSyclDeviceRef dref): cdef SyclDevice ret = SyclDevice.__new__(SyclDevice) - ret.device_ptr = dref - ret.vendor_name = DPPLDevice_GetVendorName(dref) - ret.device_name = DPPLDevice_GetName(dref) - ret.driver_version = DPPLDevice_GetDriverInfo(dref) + ret._device_ref = dref + ret._vendor_name = DPPLDevice_GetVendorName(dref) + ret._device_name = DPPLDevice_GetName(dref) + ret._driver_version = DPPLDevice_GetDriverInfo(dref) return ret def __dealloc__ (self): - DPPLDevice_Delete(self.device_ptr) - DPPLCString_Delete(self.device_name) - DPPLCString_Delete(self.vendor_name) - DPPLCString_Delete(self.driver_version) + DPPLDevice_Delete(self._device_ref) + DPPLCString_Delete(self._device_name) + DPPLCString_Delete(self._vendor_name) + DPPLCString_Delete(self._driver_version) def dump_device_info (self): ''' Print information about the SYCL device. ''' - DPPLDevice_DumpInfo(self.device_ptr) + DPPLDevice_DumpInfo(self._device_ref) def get_device_name (self): ''' Returns the name of the device as a string ''' - return self.device_name + return self._device_name.decode() def get_vendor_name (self): ''' Returns the device vendor name as a string ''' - return self.vendor_name + return self._vendor_name.decode() def get_driver_version (self): ''' Returns the OpenCL software driver version as a string @@ -103,13 +122,101 @@ cdef class SyclDevice: device is an OpenCL device. Returns a string class with the value "1.2" if this SYCL device is a host device. ''' - return self.driver_version + return self._driver_version.decode() - cdef DPPLSyclDeviceRef get_device_ptr (self): + cdef DPPLSyclDeviceRef get_device_ref (self): ''' Returns the DPPLSyclDeviceRef pointer for this class. ''' - return self.device_ptr + return self._device_ref + + +cdef class SyclEvent: + ''' Wrapper class for a Sycl Event + ''' + + @staticmethod + cdef SyclEvent _create (DPPLSyclEventRef eref, list args): + cdef SyclEvent ret = SyclEvent.__new__(SyclEvent) + ret._event_ref = eref + ret._args = args + return ret + + def __dealloc__ (self): + self.wait() + DPPLEvent_Delete(self._event_ref) + + cdef DPPLSyclEventRef get_event_ref (self): + ''' Returns the DPPLSyclEventRef pointer for this class. + ''' + return self._event_ref + + cpdef void wait (self): + DPPLEvent_Wait(self._event_ref) + + +cdef class SyclKernel: + ''' Wraps a sycl::kernel object created from an OpenCL interoperability + kernel. + ''' + + @staticmethod + cdef SyclKernel _create (DPPLSyclKernelRef kref): + cdef SyclKernel ret = SyclKernel.__new__(SyclKernel) + ret._kernel_ref = kref + ret._function_name = DPPLKernel_GetFunctionName(kref) + return ret + + def __dealloc__ (self): + DPPLKernel_Delete(self._kernel_ref) + DPPLCString_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 DPPLKernel_GetNumArgs(self._kernel_ref) + + cdef DPPLSyclKernelRef get_kernel_ref (self): + ''' Returns the DPPLSyclKernelRef pointer for this SyclKernel. + ''' + return self._kernel_ref + + +cdef class SyclProgram: + ''' Wraps a sycl::program object created from an OpenCL interoperability + program. + + SyclProgram exposes the C API from dppl_sycl_program_interface.h. A + SyclProgram can be created from either a source string or a SPIR-V + binary file. + ''' + + @staticmethod + cdef SyclProgram _create (DPPLSyclProgramRef pref): + cdef SyclProgram ret = SyclProgram.__new__(SyclProgram) + ret._program_ref = pref + return ret + + def __dealloc__ (self): + DPPLProgram_Delete(self._program_ref) + + cdef DPPLSyclProgramRef 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(DPPLProgram_GetKernel(self._program_ref, + name)) + + def has_sycl_kernel(self, str kernel_name): + name = kernel_name.encode('utf8') + return DPPLProgram_HasKernel(self._program_ref, name) + +import ctypes cdef class SyclQueue: ''' Wrapper class for a Sycl queue. @@ -118,20 +225,202 @@ cdef class SyclQueue: @staticmethod cdef SyclQueue _create (DPPLSyclQueueRef qref): cdef SyclQueue ret = SyclQueue.__new__(SyclQueue) - ret.queue_ptr = qref + ret._context = SyclContext._create(DPPLQueue_GetContext(qref)) + ret._device = SyclDevice._create(DPPLQueue_GetDevice(qref)) + ret._queue_ref = qref return ret def __dealloc__ (self): - DPPLQueue_Delete(self.queue_ptr) + DPPLQueue_Delete(self._queue_ref) + + cdef _raise_queue_submit_error (self, fname, errcode): + e = SyclKernelSubmitError("Kernel submission to Sycl queue failed.") + e.fname = fname + e.code = errcode + raise e + + cdef _raise_invalid_range_error (self, fname, ndims, errcode): + e = SyclKernelInvalidRangeError("Range with ", ndims, " not allowed. " + "Range should have between one and " + "three dimensions.") + e.fname = fname + e.code = errcode + raise e + + cdef int _populate_args (self, list args, void **kargs, \ + DPPLKernelArgType *kargty): + cdef int ret = 0 + for idx, arg in enumerate(args): + if isinstance(arg, ctypes.c_char): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._CHAR + elif isinstance(arg, ctypes.c_int): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._INT + elif isinstance(arg, ctypes.c_uint): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._UNSIGNED_INT + elif isinstance(arg, ctypes.c_long): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._LONG + elif isinstance(arg, ctypes.c_longlong): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._LONG_LONG + elif isinstance(arg, ctypes.c_ulonglong): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._UNSIGNED_LONG_LONG + elif isinstance(arg, ctypes.c_short): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._SHORT + elif isinstance(arg, ctypes.c_size_t): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._SIZE_T + elif isinstance(arg, ctypes.c_float): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._FLOAT + elif isinstance(arg, ctypes.c_double): + kargs[idx] = (ctypes.addressof(arg)) + kargty[idx] = _arg_data_type._DOUBLE + elif isinstance(arg, Memory): + kargs[idx]= (arg._pointer) + kargty[idx] = _arg_data_type._VOID_PTR + else: + ret = -1 + return ret + + + cdef int _populate_range (self, size_t Range[3], list S, size_t nS): + + cdef int ret = 0 + + if nS == 1: + Range[0] = S[0] + Range[1] = 1 + Range[2] = 1 + elif nS == 2: + Range[0] = S[0] + Range[1] = S[1] + Range[2] = 1 + elif nS == 3: + Range[0] = S[0] + Range[1] = S[1] + Range[2] = S[2] + else: + ret = -1 + + return ret + cpdef SyclContext get_sycl_context (self): - return SyclContext._create(DPPLQueue_GetContext(self.queue_ptr)) + return self._context cpdef SyclDevice get_sycl_device (self): - return SyclDevice._create(DPPLQueue_GetDevice(self.queue_ptr)) + return self._device cdef DPPLSyclQueueRef get_queue_ref (self): - return self.queue_ptr + return self._queue_ref + + cpdef SyclEvent submit (self, SyclKernel kernel, list args, list gS, \ + list lS = None, list dEvents = None): + + cdef void **kargs = NULL + cdef DPPLKernelArgType *kargty = NULL + cdef DPPLSyclEventRef *depEvents = NULL + cdef DPPLSyclEventRef Eref = NULL + cdef int ret + cdef size_t gRange[3] + cdef size_t lRange[3] + cdef size_t nGS = len(gS) + cdef size_t nLS = len(lS) if lS is not None else 0 + cdef size_t nDE = len(dEvents) if dEvents is not None else 0 + + # Allocate the arrays to be sent to DPPLQueue_Submit + kargs = malloc(len(args) * sizeof(void*)) + if not kargs: + raise MemoryError() + kargty = malloc(len(args)*sizeof(DPPLKernelArgType)) + if not kargty: + free(kargs) + raise MemoryError() + # Create the array of dependent events if any + if dEvents is not None and nDE > 0: + depEvents = malloc(nDE*sizeof(DPPLSyclEventRef)) + if not depEvents: + free(kargs) + free(kargty) + raise MemoryError() + else: + for idx, de in enumerate(dEvents): + depEvents[idx] = (de).get_event_ref() + + # populate the args and argstype arrays + ret = self._populate_args(args, kargs, kargty) + if ret == -1: + free(kargs) + free(kargty) + free(depEvents) + raise TypeError("Unsupported type for a kernel argument") + + if lS is None: + ret = self._populate_range (gRange, gS, nGS) + if ret == -1: + free(kargs) + free(kargty) + free(depEvents) + self._raise_invalid_range_error("SyclQueue.submit", nGS, -1) + + Eref = DPPLQueue_SubmitRange(kernel.get_kernel_ref(), + self.get_queue_ref(), + kargs, + kargty, + len(args), + gRange, + nGS, + depEvents, + nDE) + else: + ret = self._populate_range (gRange, gS, nGS) + if ret == -1: + free(kargs) + free(kargty) + free(depEvents) + self._raise_invalid_range_error("SyclQueue.submit", nGS, -1) + ret = self._populate_range (lRange, lS, nLS) + if ret == -1: + free(kargs) + free(kargty) + free(depEvents) + self._raise_invalid_range_error("SyclQueue.submit", nLS, -1) + + if nGS != nLS: + free(kargs) + free(kargty) + free(depEvents) + raise ValueError("Local and global ranges need to have same " + "number of dimensions.") + + Eref = DPPLQueue_SubmitNDRange(kernel.get_kernel_ref(), + self.get_queue_ref(), + kargs, + kargty, + len(args), + gRange, + lRange, + nGS, + depEvents, + nDE) + free(kargs) + free(kargty) + free(depEvents) + + if Eref is NULL: + # \todo get the error number from dpctl-capi + self._raise_queue_submit_error("DPPLQueue_Submit", -1) + + return SyclEvent._create(Eref, args) + + cpdef void wait (self): + DPPLQueue_Wait(self._queue_ref) cpdef memcpy (self, dest, src, int count): cdef void *c_dest @@ -147,21 +436,24 @@ cdef class SyclQueue: else: raise TypeError("Parameter src should be Memory.") - DPPLQueue_Memcpy(self.queue_ptr, c_dest, c_src, count) + DPPLQueue_Memcpy(self._queue_ref, c_dest, c_src, count) cdef class _SyclQueueManager: + ''' Wrapper for the C API's sycl queue manager interface. + ''' + def _set_as_current_queue (self, device_ty, device_id): - cdef DPPLSyclQueueRef queue_ptr + cdef DPPLSyclQueueRef queue_ref if device_ty == device_type.gpu: - queue_ptr = DPPLQueueMgr_PushQueue(_device_type._GPU, device_id) + queue_ref = DPPLQueueMgr_PushQueue(_device_type._GPU, device_id) elif device_ty == device_type.cpu: - queue_ptr = DPPLQueueMgr_PushQueue(_device_type._CPU, device_id) + queue_ref = DPPLQueueMgr_PushQueue(_device_type._CPU, device_id) else: e = UnsupportedDeviceTypeError("Device can only be cpu or gpu") raise e - return SyclQueue._create(queue_ptr) + return SyclQueue._create(queue_ref) def _remove_current_queue (self): DPPLQueueMgr_PopQueue() @@ -223,6 +515,7 @@ cdef class _SyclQueueManager: else: return False + # This private instance of the _SyclQueueManager should not be directly # accessed outside the module. _qmgr = _SyclQueueManager() @@ -238,6 +531,69 @@ has_sycl_platforms = _qmgr.has_sycl_platforms set_default_queue = _qmgr.set_default_queue is_in_device_context = _qmgr.is_in_device_context + +def create_program_from_source (SyclQueue q, unicode source, unicode copts=""): + ''' Creates a Sycl interoperability program from an OpenCL source string. + + We use the DPPLProgram_CreateFromOCLSource() C API function to create + a Sycl progrma from an OpenCL source program that can contain multiple + kernels. + + Parameters: + q (SyclQueue) : The SyclQueue object wraps the Sycl device for + which the program will 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 SyclProgram object wrapping the + syc::program returned by the C API. + ''' + + cdef DPPLSyclProgramRef Pref + + cdef bytes bSrc = source.encode('utf8') + cdef bytes bCOpts = copts.encode('utf8') + cdef const char *Src = bSrc + cdef const char *COpts = bCOpts + cdef DPPLSyclContextRef CRef = q.get_sycl_context().get_context_ref() + Pref = DPPLProgram_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 DPPLProgram_CreateFromOCLSpirv() C API function to create + a Sycl progrma from an compiled SPIR-V binary file. + + Parameters: + q (SyclQueue): The SyclQueue object wraps the Sycl device for + which the program will be built. + IL (const char[:]) : SPIR-V binary IL file for an OpenCL program. + + Returns: + program (SyclProgram): A SyclProgram object wrapping the + syc::program returned by the C API. + ''' + + cdef DPPLSyclProgramRef Pref + cdef const unsigned char *dIL = &IL[0] + cdef DPPLSyclContextRef CRef = q.get_sycl_context().get_context_ref() + cdef size_t length = IL.shape[0] + Pref = DPPLProgram_CreateFromOCLSpirv(CRef, dIL, length) + if Pref is NULL: + raise SyclProgramCompilationError() + + return SyclProgram._create(Pref) + + from contextlib import contextmanager @contextmanager diff --git a/dpctl/tests/__init__.py b/dpctl/tests/__init__.py index ce80b2561c..52c61b29d1 100644 --- a/dpctl/tests/__init__.py +++ b/dpctl/tests/__init__.py @@ -22,6 +22,9 @@ ## Top-level module of all dpctl Python unit test cases. ##===----------------------------------------------------------------------===## +from .test_dump_functions import * +from .test_sycl_kernel_submit import * +from .test_sycl_program import * from .test_sycl_queue_manager import * +from .test_sycl_queue_memcpy import * from .test_sycl_usm import * -from .test_dump_functions import * diff --git a/dpctl/tests/input_files/multi_kernel.spv b/dpctl/tests/input_files/multi_kernel.spv new file mode 100644 index 0000000000..1fa772d8df Binary files /dev/null and b/dpctl/tests/input_files/multi_kernel.spv differ diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py new file mode 100644 index 0000000000..69b24c7405 --- /dev/null +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -0,0 +1,69 @@ +##===------------- test_sycl_kernel_submit.py - dpctl -----*- Python -*---===## +## +## 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 +## Defines unit test cases for kernel submission to a sycl::queue. +## +##===----------------------------------------------------------------------===## +import ctypes +import dpctl +import unittest +import dpctl._memory as dpctl_mem +import numpy as np + +@unittest.skipIf(not dpctl.has_sycl_platforms(), "No SYCL platforms available") +class Test1DKernelSubmit (unittest.TestCase): + + def test_create_program_from_source (self): + oclSrc = " \ + kernel void axpy(global int* a, global int* b, global int* c, int d) { \ + size_t index = get_global_id(0); \ + c[index] = d*a[index] + b[index]; \ + }" + with dpctl.device_context(dpctl.device_type.gpu, 0): + q = dpctl.get_current_queue() + prog = dpctl.create_program_from_source(q, oclSrc) + axpyKernel = prog.get_sycl_kernel('axpy') + + abuf = dpctl_mem.MemoryUSMShared(1024*np.dtype('i').itemsize) + bbuf = dpctl_mem.MemoryUSMShared(1024*np.dtype('i').itemsize) + cbuf = dpctl_mem.MemoryUSMShared(1024*np.dtype('i').itemsize) + a = np.ndarray((1024), buffer=abuf, dtype='i') + b = np.ndarray((1024), buffer=bbuf, dtype='i') + c = np.ndarray((1024), buffer=cbuf, dtype='i') + a[:] = np.arange(1024) + b[:] = np.arange(1024, 0, -1) + c[:] = 0 + d = 2 + args = [] + + args.append(a.base) + args.append(b.base) + args.append(c.base) + args.append(ctypes.c_int(d)) + + r = [1024] + + q.submit(axpyKernel, args, r) + self.assertTrue(np.allclose(c, a*d + b)) + + +if __name__ == '__main__': + unittest.main() diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py new file mode 100644 index 0000000000..1a87de88d3 --- /dev/null +++ b/dpctl/tests/test_sycl_program.py @@ -0,0 +1,87 @@ +##===------------- test_sycl_program.py - dpctl -------*- Python -*-------===## +## +## 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 +## Defines unit test cases for the SyclProgram and SyclKernel classes defined +## in sycl_core.pyx. +##===----------------------------------------------------------------------===## + +import dpctl +import unittest +import os + + +@unittest.skipIf(not dpctl.has_gpu_queues(), "No SYCL platforms available") +class TestProgramFromOCLSource (unittest.TestCase): + + def test_create_program_from_source (self): + oclSrc = " \ + kernel void add(global int* a, global int* b, global int* c) { \ + size_t index = get_global_id(0); \ + c[index] = a[index] + b[index]; \ + } \ + kernel void axpy(global int* a, global int* b, global int* c, int d) { \ + size_t index = get_global_id(0); \ + c[index] = a[index] + d*b[index]; \ + }" + with dpctl.device_context(dpctl.device_type.gpu, 0): + q = dpctl.get_current_queue() + prog = dpctl.create_program_from_source(q, oclSrc) + self.assertIsNotNone(prog) + + self.assertTrue(prog.has_sycl_kernel("add")) + self.assertTrue(prog.has_sycl_kernel("axpy")) + + addKernel = prog.get_sycl_kernel('add') + axpyKernel = prog.get_sycl_kernel('axpy') + + self.assertEqual(addKernel.get_function_name(),"add") + self.assertEqual(axpyKernel.get_function_name(),"axpy") + self.assertEqual(addKernel.get_num_args(), 3) + self.assertEqual(axpyKernel.get_num_args(), 4) + + +@unittest.skipIf(not dpctl.has_gpu_queues(), "No SYCL platforms available") +class TestProgramFromSPRIV (unittest.TestCase): + + def test_create_program_from_spirv(self): + + CURR_DIR = os.path.dirname(os.path.abspath(__file__)) + spirv_file = os.path.join(CURR_DIR, 'input_files/multi_kernel.spv') + with open(spirv_file, 'rb') as fin: + spirv = fin.read() + with dpctl.device_context(dpctl.device_type.gpu, 0): + q = dpctl.get_current_queue() + prog = dpctl.create_program_from_spirv(q,spirv) + self.assertIsNotNone(prog) + self.assertTrue(prog.has_sycl_kernel("add")) + self.assertTrue(prog.has_sycl_kernel("axpy")) + + addKernel = prog.get_sycl_kernel('add') + axpyKernel = prog.get_sycl_kernel('axpy') + + self.assertEqual(addKernel.get_function_name(),"add") + self.assertEqual(axpyKernel.get_function_name(),"axpy") + self.assertEqual(addKernel.get_num_args(), 3) + self.assertEqual(axpyKernel.get_num_args(), 4) + + +if __name__ == '__main__': + unittest.main() diff --git a/scripts/build_for_develop.sh b/scripts/build_for_develop.sh index 1f673b5bc4..63057d8099 100755 --- a/scripts/build_for_develop.sh +++ b/scripts/build_for_develop.sh @@ -22,10 +22,11 @@ cmake \ -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ -DPYTHON_INCLUDE_DIR=${PYTHON_INC} \ -DNUMPY_INCLUDE_DIR=${NUMPY_INC} \ + -DGTEST_INCLUDE_DIR=${CONDA_PREFIX}/include/ \ + -DGTEST_LIB_DIR=${CONDA_PREFIX}/lib \ ../backends -make V=1 -n -j 4 && make install -#make check +make V=1 -n -j 4 && make check && make install popd cp install/lib/*.so dpctl/ @@ -46,3 +47,4 @@ export CXX=dpcpp export CFLAGS=-fPIC python setup.py clean --all python setup.py build develop +python -m unittest dpctl.tests