From 5853eda9609c1b5116344c8b02581c120bfc6429 Mon Sep 17 00:00:00 2001 From: vlad-perevezentsev <68376232+vlad-perevezentsev@users.noreply.github.com> Date: Tue, 22 Dec 2020 21:48:52 +0300 Subject: [PATCH 1/8] Fix incorrect import (#234) --- examples/cython/usm_memory/run.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/cython/usm_memory/run.py b/examples/cython/usm_memory/run.py index 422c4baaf1..0373c25217 100644 --- a/examples/cython/usm_memory/run.py +++ b/examples/cython/usm_memory/run.py @@ -1,5 +1,5 @@ # coding: utf-8 -import dpctl._memory as dpctl_mem +import dpctl.memory as dpctl_mem import blackscholes_usm as bs import numpy as np, dpctl from reference_black_scholes import ref_python_black_scholes From fdb5569a8b81fdc389e47bef8f6f7ca101d23697 Mon Sep 17 00:00:00 2001 From: Diptorup Deb <3046810+diptorupd@users.noreply.github.com> Date: Wed, 20 Jan 2021 12:01:43 -0600 Subject: [PATCH 2/8] Cmake improvements and Coverage for C API (#242) * Cmake build script changes for dpCtl C API. -- Create cmake modules to store logic to search for dpcpp, level zero, llvm-cov, lcov. -- Add new Cmake options to control level zero program creation and code coverage reporting. -- Two new rules llvm-cov and lcov-genhtml to produce code coverage metrics. --- dpctl-capi/CMakeLists.txt | 163 ++++++++++--------- dpctl-capi/cmake/modules/FindDPCPP.cmake | 115 +++++++++++++ dpctl-capi/cmake/modules/FindLcov.cmake | 33 ++++ dpctl-capi/cmake/modules/FindLevelZero.cmake | 34 ++++ dpctl-capi/dbg_build.sh | 20 ++- dpctl-capi/include/Config/.gitignore | 1 + dpctl-capi/include/Config/dpctl_config.h.in | 32 ++++ dpctl-capi/tests/CMakeLists.txt | 50 +++++- scripts/build_backend.py | 6 +- scripts/build_for_develop.sh | 12 +- 10 files changed, 375 insertions(+), 91 deletions(-) create mode 100644 dpctl-capi/cmake/modules/FindDPCPP.cmake create mode 100644 dpctl-capi/cmake/modules/FindLcov.cmake create mode 100644 dpctl-capi/cmake/modules/FindLevelZero.cmake create mode 100644 dpctl-capi/include/Config/.gitignore create mode 100644 dpctl-capi/include/Config/dpctl_config.h.in diff --git a/dpctl-capi/CMakeLists.txt b/dpctl-capi/CMakeLists.txt index b327f271c9..16748f66cf 100644 --- a/dpctl-capi/CMakeLists.txt +++ b/dpctl-capi/CMakeLists.txt @@ -1,45 +1,39 @@ cmake_minimum_required(VERSION 3.10 FATAL_ERROR) -project("dpCtl - A lightweight SYCL wrapper for Python") +project("dpCtl C API - A C wrapper for a subset of SYCL") -# The function checks is DPCPP_ROOT is valid and points to a dpcpp installation -function (check_for_dpcpp) - string(COMPARE EQUAL "${DPCPP_ROOT}" "" no_dpcpp_root) - if(${no_dpcpp_root}) - message(FATAL_ERROR "Set the DPCPP_ROOT argument providing the path to \ - a dpcpp installation.") - endif() - - if(WIN32) - set (dpcpp_cmd "${DPCPP_ROOT}/bin/dpcpp") - set (dpcpp_arg "--version") - elseif(UNIX) - set (dpcpp_cmd "${DPCPP_ROOT}/bin/dpcpp") - set (dpcpp_arg "--version") - else() - message(FATAL_ERROR "Unsupported system.") - endif() +# Option to turn on support for creating Level Zero interoperability programs +# from a SPIR-V binary file. +option(DPCTL_ENABLE_LO_PROGRAM_CREATION + "Enable Level Zero Program creation from SPIR-V" + OFF +) +# Option to generate code coverage report using llvm-cov and lcov. +option(DPCTL_GENERATE_COVERAGE + "Build dpctl C API with coverage instrumentation instrumentation" + OFF +) +# Option to output html coverage report at a specific location. +option(DPCTL_COVERAGE_REPORT_OUTPUT_DIR + "Save the generated lcov html report to the specified location" + OFF +) +# Option to build the Gtests for dpctl C API +option(DPCTL_BUILD_CAPI_TESTS + "Build dpctl C API google tests" + OFF +) - # Check if dpcpp is available - execute_process( - COMMAND ${dpcpp_cmd} ${dpcpp_arg} - WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} - RESULT_VARIABLE dpcpp_result - OUTPUT_VARIABLE dpcpp_ver - ) +# Load our CMake modules to search for DPCPP and Level Zero +set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/modules/") +find_package(DPCPP REQUIRED) - if(${dpcpp_result} MATCHES "0") - string(REPLACE "\n" ";" DPCPP_VERSION_LIST "${dpcpp_ver}") - list(GET DPCPP_VERSION_LIST 0 dpcpp_ver_line) - foreach(X ${DPCPP_VERSION_LIST}) - message(STATUS "dpcpp ver[${dpcpp_result}]: ${X}") - endforeach() - else() - message(FATAL_ERROR "DPCPP needed to build dpctl_sycl_interface") - endif() -endfunction() +if(DPCTL_ENABLE_LO_PROGRAM_CREATION) + set(DPCTL_ENABLE_LO_PROGRAM_CREATION 1) + find_package(LevelZero REQUIRED) +endif() -# Check for dpcpp in the specified DPCPP_ROOT -check_for_dpcpp() +configure_file(${CMAKE_SOURCE_DIR}/include/Config/dpctl_config.h.in + ${CMAKE_SOURCE_DIR}/include/Config/dpctl_config.h) if(WIN32) set(CMAKE_CXX_COMPILER:PATH "${DPCPP_ROOT}/bin/dpcpp") @@ -54,6 +48,9 @@ if(WIN32) set(CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG} ${WARNING_FLAGS} -ggdb3 -DDEBUG") set(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} ${WARNING_FLAGS} -ggdb3 -DDEBUG -Qstd=c++17") elseif(UNIX) + set(CMAKE_CXX_COMPILER:PATH "${DPCPP_ROOT}/bin/dpcpp") + set(CMAKE_C_COMPILER:PATH "${DPCPP_ROOT}/bin/clang") + set(CMAKE_LINKER:PATH "${DPCPP_ROOT}/bin/lld") set(SDL_FLAGS "-fstack-protector -fstack-protector-all -fpic -fPIC -D_FORTIFY_SOURCE=2 -Wformat -Wformat-security -fno-strict-overflow -fno-delete-null-pointer-checks") set(WARNING_FLAGS "-Wall -Wextra -Winit-self -Wunused-function -Wuninitialized -Wmissing-declarations -fdiagnostics-color=auto") set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${WARNING_FLAGS} ${SDL_FLAGS}") @@ -80,30 +77,32 @@ add_library( helper/source/dpctl_utils_helper.cpp ) -# Install DPCTLSyclInterface -target_include_directories( - DPCTLSyclInterface - PRIVATE - ${CMAKE_SOURCE_DIR}/include/ - ${CMAKE_SOURCE_DIR}/helper/include/ +target_include_directories(DPCTLSyclInterface + PRIVATE + ${CMAKE_SOURCE_DIR}/include/ + ${CMAKE_SOURCE_DIR}/helper/include/ + ${DPCPP_SYCL_INCLUDE_DIR} +) +target_link_libraries(DPCTLSyclInterface + PRIVATE ${DPCPP_SYCL_LIBRARY} + PRIVATE ${DPCPP_OPENCL_LIBRARY} ) -if(WIN32) - message( - STATUS - "SYCL_INCLUDE_DIR: " - ${DPCPP_ROOT}/include/sycl - ) - target_include_directories( - DPCTLSyclInterface - PUBLIC - ${DPCPP_ROOT}/include/sycl - ) - target_link_libraries( - DPCTLSyclInterface - PRIVATE ${DPCPP_ROOT}/lib/sycl.lib - PRIVATE ${DPCPP_ROOT}/lib/OpenCL.lib - ) +if(DPCTL_ENABLE_LO_PROGRAM_CREATION) + if(UNIX) + target_include_directories(DPCTLSyclInterface + PRIVATE + ${LEVEL_ZERO_INCLUDE_DIR} + ) + target_link_libraries(DPCTLSyclInterface + PRIVATE ${LEVEL_ZERO_LIBRARY} + ) + else() + message(WARNING + "DPCTL support Level Zero program creation not supported " + "on this system." + ) + endif() endif() install( @@ -114,34 +113,52 @@ install( ) # Install all headers -file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/include/*.h*") +file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/include/*.h") foreach(HEADER ${HEADERS}) install(FILES "${HEADER}" DESTINATION include) endforeach() # Install all headers in include/Support -file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/include/Support/*.h*") +file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/include/Support/*.h") foreach(HEADER ${HEADERS}) install(FILES "${HEADER}" DESTINATION include/Support) endforeach() # Install all headers in helper/include -file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/helper/include/*.h*") +file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/helper/include/*.h") foreach(HEADER ${HEADERS}) install(FILES "${HEADER}" DESTINATION helper/include) endforeach() -option( - BUILD_CAPI_TESTS - "Build dpctl C API google tests" - OFF -) +# Install all headers in include/Config +file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/include/Config/*.h") +foreach(HEADER ${HEADERS}) + install(FILES "${HEADER}" DESTINATION include/Config) +endforeach() -# Enable to build the dpCtl backend test cases -if(BUILD_CAPI_TESTS) - add_subdirectory(tests) +# Enable code coverage related settings +if(DPCTL_GENERATE_COVERAGE) + # check if llvm-cov and lcov are available + find_package(Lcov REQUIRED) + # These flags are set inside FindDPCPP + if(NOT (${LLVM_COV_FOUND} AND ${LLVM_PROFDATA_FOUND})) + message(FATAL_ERROR + "llvm-cov and llvm-profdata are needed to generate coverage." + ) + endif() + # Turn on DPCTL_BUILD_CAPI_TESTS as building tests is needed to generate + # coverage reports + set(DPCTL_BUILD_CAPI_TESTS "ON") + if(DPCTL_COVERAGE_REPORT_OUTPUT_DIR) + set(COVERAGE_OUTPUT_DIR ${DPCTL_COVERAGE_REPORT_OUTPUT_DIR}) + message(STATUS "Coverage reports to be saved at ${COVERAGE_OUTPUT_DIR}") + else() + set(COVERAGE_OUTPUT_DIR ${CMAKE_CURRENT_BINARY_DIR}) + message(STATUS "Coverage reports to be saved at ${COVERAGE_OUTPUT_DIR}") + endif() endif() - -# Todo : Add build rules for doxygen -# maybe refer https://devblogs.microsoft.com/cppblog/clear-functional-c-documentation-with-sphinx-breathe-doxygen-cmake/ +# Add sub-directory to build the dpCtl C API test cases +if(DPCTL_BUILD_CAPI_TESTS) + add_subdirectory(tests) +endif() diff --git a/dpctl-capi/cmake/modules/FindDPCPP.cmake b/dpctl-capi/cmake/modules/FindDPCPP.cmake new file mode 100644 index 0000000000..370dd965af --- /dev/null +++ b/dpctl-capi/cmake/modules/FindDPCPP.cmake @@ -0,0 +1,115 @@ +# Data Parallel Control Library (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. +# +# CMake find_package() module for the DPCPP compiler and development +# environment. +# +# Example usage: +# +# find_package(DPCPP) +# +# If successful, the following variables will be defined: +# DPCPP_FOUND +# DPCPP_VERSION +# DPCPP_INCLUDE_DIR +# DPCPP_SYCL_INCLUDE_DIR +# DPCPP_LIBRARY_DIR +# DPCPP_SYCL_LIBRARY +# DPCPP_OPENCL_LIBRARY + +include( FindPackageHandleStandardArgs ) + +string(COMPARE EQUAL "${DPCPP_INSTALL_DIR}" "" no_dpcpp_root) +if(${no_dpcpp_root}) + message(STATUS "Set the DPCPP_ROOT argument providing the path to \ + a dpcpp installation.") + return() +endif() + +if(WIN32 OR UNIX) + set(dpcpp_cmd "${DPCPP_INSTALL_DIR}/bin/dpcpp") + set(dpcpp_arg "--version") +else() + message(FATAL_ERROR "Unsupported system.") +endif() + +# Check if dpcpp is available +execute_process( + COMMAND ${dpcpp_cmd} ${dpcpp_arg} + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} + RESULT_VARIABLE dpcpp_result + OUTPUT_VARIABLE dpcpp_ver +) + +# If dpcpp is found then set then set the package variables +if(${dpcpp_result} MATCHES "0") + string(REPLACE "\n" ";" DPCPP_VERSION_LIST "${dpcpp_ver}") + list(GET DPCPP_VERSION_LIST 0 dpcpp_ver_line) + foreach(X ${DPCPP_VERSION_LIST}) + message(STATUS "dpcpp ver[${dpcpp_result}]: ${X}") + endforeach() + + # check if llvm-cov and llvm-profdata are packaged as part of dpcpp + find_program(LLVM_COV_EXE + llvm-cov + PATHS ${DPCPP_INSTALL_DIR}/bin + NO_DEFAULT_PATH + ) + + if(LLVM_COV_EXE) + set(LLVM_COV_FOUND TRUE) + else() + set(LLVM_COV_FOUND FALSE) + endif() + + find_program(LLVM_PROFDATA_EXE + llvm-profdata + PATHS ${DPCPP_INSTALL_DIR}/bin + NO_DEFAULT_PATH + ) + + if(LLVM_PROFDATA_EXE) + set(LLVM_PROFDATA_FOUND TRUE) + else() + set(LLVM_PROFDATA_FOUND FALSE) + endif() + + # set package-level variables + set(DPCPP_ROOT ${DPCPP_INSTALL_DIR}) + list(POP_FRONT DPCPP_VERSION_LIST DPCPP_VERSION) + set(DPCPP_INCLUDE_DIR ${DPCPP_INSTALL_DIR}/include) + set(DPCPP_SYCL_INCLUDE_DIR ${DPCPP_INSTALL_DIR}/include/sycl) + set(DPCPP_LIBRARY_DIR ${DPCPP_INSTALL_DIR}/lib) + if(WIN32) + set(DPCPP_SYCL_LIBRARY ${DPCPP_INSTALL_DIR}/lib/sycl.lib) + set(DPCPP_OPENCL_LIBRARY ${DPCPP_INSTALL_DIR}/lib/OpenCL.lib) + elseif(UNIX) + set(DPCPP_SYCL_LIBRARY ${DPCPP_INSTALL_DIR}/lib/libsycl.so) + set(DPCPP_OPENCL_LIBRARY ${DPCPP_INSTALL_DIR}/lib/libOpenCL.so) + endif() +else() + message(STATUS "DPCPP needed to build dpctl_sycl_interface") + return() +endif() + +find_package_handle_standard_args(DPCPP DEFAULT_MSG + DPCPP_VERSION + DPCPP_INCLUDE_DIR + DPCPP_SYCL_INCLUDE_DIR + DPCPP_LIBRARY_DIR + DPCPP_SYCL_LIBRARY + DPCPP_OPENCL_LIBRARY +) diff --git a/dpctl-capi/cmake/modules/FindLcov.cmake b/dpctl-capi/cmake/modules/FindLcov.cmake new file mode 100644 index 0000000000..43666b2618 --- /dev/null +++ b/dpctl-capi/cmake/modules/FindLcov.cmake @@ -0,0 +1,33 @@ +# Data Parallel Control Library (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. +# +# CMake find_package() module for lcov. +# +# Example usage: +# +# find_package(Lcov) +# +# If successful, the following variables will be defined: +# LCOV_EXE- The path to lcov executable +# LCOV_FOUND + +find_program(LCOV_EXE lcov) +find_program(GENHTML_EXE genhtml) + +find_package_handle_standard_args(Lcov DEFAULT_MSG + LCOV_EXE + GENHTML_EXE +) diff --git a/dpctl-capi/cmake/modules/FindLevelZero.cmake b/dpctl-capi/cmake/modules/FindLevelZero.cmake new file mode 100644 index 0000000000..f58e9e7a2b --- /dev/null +++ b/dpctl-capi/cmake/modules/FindLevelZero.cmake @@ -0,0 +1,34 @@ +# Data Parallel Control Library (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. +# +# CMake find_package() module for the Level Zero loader library and headers. +# +# Example usage: +# +# find_package(LevelZero) +# +# If successful, the following variables will be defined: +# LEVEL_ZERO_INCLUDE_DIR +# LEVEL_ZERO_LIBRARY - the full path to the ze_loader library +# TODO: Add a way to record the version of the level_zero library + +find_library(LEVEL_ZERO_LIBRARY ze_loader) +find_path(LEVEL_ZERO_INCLUDE_DIR NAMES level_zero/zet_api.h) + +find_package_handle_standard_args(LevelZero DEFAULT_MSG + LEVEL_ZERO_INCLUDE_DIR + LEVEL_ZERO_LIBRARY +) diff --git a/dpctl-capi/dbg_build.sh b/dpctl-capi/dbg_build.sh index 953d9d275c..60c8674a42 100755 --- a/dpctl-capi/dbg_build.sh +++ b/dpctl-capi/dbg_build.sh @@ -9,18 +9,22 @@ rm -rf ${INSTALL_PREFIX} export ONEAPI_ROOT=/opt/intel/oneapi DPCPP_ROOT=${ONEAPI_ROOT}/compiler/latest/linux -cmake \ - -DCMAKE_BUILD_TYPE=Debug \ - -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ - -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ - -DDPCPP_ROOT=${DPCPP_ROOT} \ - -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ - -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ - -DBUILD_CAPI_TESTS=ON \ +cmake \ + -DCMAKE_BUILD_TYPE=Debug \ + -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ + -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ + -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ + -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ + -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ + -DDPCTL_BUILD_CAPI_TESTS=ON \ + -DDPCTL_GENERATE_COVERAGE=ON \ .. make V=1 -n -j 4 && make check && make install +# Turn on to generate coverage report html files +make lcov-genhtml + # For more verbose tests use: # cd tests # ctest -V --progress --output-on-failure -j 4 diff --git a/dpctl-capi/include/Config/.gitignore b/dpctl-capi/include/Config/.gitignore new file mode 100644 index 0000000000..ecaeb22660 --- /dev/null +++ b/dpctl-capi/include/Config/.gitignore @@ -0,0 +1 @@ +dpctl_config.h diff --git a/dpctl-capi/include/Config/dpctl_config.h.in b/dpctl-capi/include/Config/dpctl_config.h.in new file mode 100644 index 0000000000..e1e3442fda --- /dev/null +++ b/dpctl-capi/include/Config/dpctl_config.h.in @@ -0,0 +1,32 @@ +//===---- dpctl-capi/Config/dpCtl-config.h - dpctl-C API -------*- C++ -*-===// +// +// Data Parallel Control Library (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 exports a set of dpCtl C API configurations. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +/* Defined when dpCtl was built with level zero program creation enabled. */ +#cmakedefine DPCTL_ENABLE_LO_PROGRAM_CREATION @DPCTL_ENABLE_LO_PROGRAM_CREATION@ + +/* The DPCPP version used to build dpCtl */ +#define DPCTL_DPCPP_VERSION "@DPCPP_VERSION@" diff --git a/dpctl-capi/tests/CMakeLists.txt b/dpctl-capi/tests/CMakeLists.txt index 8c599ce333..87dd681743 100644 --- a/dpctl-capi/tests/CMakeLists.txt +++ b/dpctl-capi/tests/CMakeLists.txt @@ -18,10 +18,50 @@ foreach(tf ${spirv-test-files}) file(COPY ${tf} DESTINATION ${CMAKE_CURRENT_BINARY_DIR}) endforeach() -file(GLOB_RECURSE sources ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp) -add_executable(dpctl_c_api_tests EXCLUDE_FROM_ALL ${sources}) -target_link_libraries( - dpctl_c_api_tests ${CMAKE_THREAD_LIBS_INIT} GTest::GTest DPCTLSyclInterface -) +if(DPCTL_GENERATE_COVERAGE) + file(GLOB_RECURSE sources ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp) + file(GLOB_RECURSE dpctl_sources ${CMAKE_CURRENT_SOURCE_DIR}/../source/*.cpp) + + # Add profiling flags + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fprofile-instr-generate -fcoverage-mapping") + + # Add all dpctl sources into a single executable so that we can run coverage + # analysis and generate a report. + add_executable(dpctl_c_api_tests + EXCLUDE_FROM_ALL + ${sources} + ${dpctl_sources} + ) + target_link_libraries(dpctl_c_api_tests + ${CMAKE_THREAD_LIBS_INIT} + GTest::GTest + DPCTLSyclInterface + ${LEVEL_ZERO_LIBRARY} + ${DPCPP_OPENCL_LIBRARY} + ) + add_custom_target(llvm-cov + COMMAND ${CMAKE_MAKE_PROGRAM} dpctl_c_api_tests + COMMAND ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests + COMMAND ${LLVM_PROFDATA_EXE} merge -sparse default.profraw -o dpctl.profdata + COMMAND ${LLVM_COV_EXE} report ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests -instr-profile=dpctl.profdata ${dpctl_sources} + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} + ) + add_custom_target(lcov-genhtml + COMMAND ${CMAKE_MAKE_PROGRAM} llvm-cov + COMMAND ${LLVM_COV_EXE} export -format=lcov ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests -instr-profile=dpctl.profdata ${dpctl_sources} > dpctl.lcov + COMMAND ${GENHTML_EXE} ${CMAKE_CURRENT_BINARY_DIR}/dpctl.lcov --output-directory ${COVERAGE_OUTPUT_DIR}/dpctl-c-api-coverage + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} + ) +else() + file(GLOB_RECURSE sources ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp) + add_executable(dpctl_c_api_tests EXCLUDE_FROM_ALL ${sources}) + target_link_libraries(dpctl_c_api_tests + ${CMAKE_THREAD_LIBS_INIT} + GTest::GTest + DPCTLSyclInterface + ${LEVEL_ZERO_LIBRARY} + ) +endif() + gtest_discover_tests(dpctl_c_api_tests) add_dependencies(check dpctl_c_api_tests) diff --git a/scripts/build_backend.py b/scripts/build_backend.py index ac61a52b98..2708094dd6 100644 --- a/scripts/build_backend.py +++ b/scripts/build_backend.py @@ -40,13 +40,13 @@ "-DCMAKE_BUILD_TYPE=Release", "-DCMAKE_INSTALL_PREFIX=" + INSTALL_PREFIX, "-DCMAKE_PREFIX_PATH=" + INSTALL_PREFIX, - "-DDPCPP_ROOT=" + DPCPP_ROOT, + "-DDPCPP_INSTALL_DIR=" + DPCPP_ROOT, "-DCMAKE_C_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang"), "-DCMAKE_CXX_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang++"), backends, ] subprocess.check_call(cmake_args, stderr=subprocess.STDOUT, shell=False) - subprocess.check_call(["make", "-j", "4"]) + subprocess.check_call(["make", "V=1", "-j", "4"]) subprocess.check_call(["make", "install"]) os.chdir(dpctl_dir) @@ -61,7 +61,7 @@ "-DCMAKE_BUILD_TYPE=Release", "-DCMAKE_INSTALL_PREFIX=" + INSTALL_PREFIX, "-DCMAKE_PREFIX_PATH=" + INSTALL_PREFIX, - "-DDPCPP_ROOT=" + DPCPP_ROOT, + "-DDPCPP_INSTALL_DIR=" + DPCPP_ROOT, backends, ] subprocess.check_call(cmake_args, stderr=subprocess.STDOUT, shell=True) diff --git a/scripts/build_for_develop.sh b/scripts/build_for_develop.sh index 0154782c9c..bd10d01e25 100755 --- a/scripts/build_for_develop.sh +++ b/scripts/build_for_develop.sh @@ -14,14 +14,22 @@ cmake \ -DCMAKE_BUILD_TYPE=Debug \ -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ - -DDPCPP_ROOT=${DPCPP_ROOT} \ + -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ - -DBUILD_CAPI_TESTS=ON \ + -DDPCTL_BUILD_CAPI_TESTS=ON \ + -DDPCTL_GENERATE_COVERAGE=ON \ ../dpctl-capi make V=1 -n -j 4 && make check && make install +if [ $? -ne 0 ]; then + echo "Building of libDPCTLSyclInterface failed. Abort!" + exit 1 +fi + +# To run code coverage for dpctl-c API +make llvm-cov # For more verbose tests use: # cd tests # ctest -V --progress --output-on-failure -j 4 From 91ef1f579a242e2dc170d50f52a93c582d260e1c Mon Sep 17 00:00:00 2001 From: Sergey Pokhodenko Date: Fri, 22 Jan 2021 15:17:28 +0300 Subject: [PATCH 3/8] Update README.md --- docs/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/README.md b/docs/README.md index fe84f32af7..4dfa01046b 100644 --- a/docs/README.md +++ b/docs/README.md @@ -16,7 +16,7 @@ documents in the current source directory in a sub-directory called `generated_docs`. The `make Sphinx` command will generate standalone Doxygen documentation and -a consolidated Sphix documentation for both dpCtl Python and C APIs. +a consolidated Sphinx documentation for both dpCtl Python and C APIs. Prerequisite ============ From 7a3d1b2fb52df6e4b76980df2548268223392c61 Mon Sep 17 00:00:00 2001 From: Diptorup Deb <3046810+diptorupd@users.noreply.github.com> Date: Fri, 22 Jan 2021 09:39:07 -0600 Subject: [PATCH 4/8] Level zero codegen (#241) * Add support to build a Level Zero interoperability SYCL program from SPIR-V binary on Linux. --- dpctl-capi/dbg_build.sh | 19 ++--- .../include/dpctl_sycl_program_interface.h | 9 ++- .../source/dpctl_sycl_program_interface.cpp | 77 ++++++++++++++++--- .../tests/test_sycl_program_interface.cpp | 46 ++++++++--- dpctl/_backend.pxd | 9 ++- dpctl/_sycl_core.pyx | 4 +- dpctl/program/_program.pxd | 3 +- dpctl/program/_program.pyx | 9 ++- dpctl/tests/test_sycl_program.py | 11 ++- scripts/build_backend.py | 1 + 10 files changed, 146 insertions(+), 42 deletions(-) diff --git a/dpctl-capi/dbg_build.sh b/dpctl-capi/dbg_build.sh index 60c8674a42..5b42f82273 100755 --- a/dpctl-capi/dbg_build.sh +++ b/dpctl-capi/dbg_build.sh @@ -9,15 +9,16 @@ rm -rf ${INSTALL_PREFIX} export ONEAPI_ROOT=/opt/intel/oneapi DPCPP_ROOT=${ONEAPI_ROOT}/compiler/latest/linux -cmake \ - -DCMAKE_BUILD_TYPE=Debug \ - -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ - -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ - -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ - -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ - -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ - -DDPCTL_BUILD_CAPI_TESTS=ON \ - -DDPCTL_GENERATE_COVERAGE=ON \ +cmake \ + -DCMAKE_BUILD_TYPE=Debug \ + -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ + -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ + -DDPCPP_INSTALL_DIR=${DPCPP_ROOT} \ + -DCMAKE_C_COMPILER:PATH=${DPCPP_ROOT}/bin/clang \ + -DCMAKE_CXX_COMPILER:PATH=${DPCPP_ROOT}/bin/dpcpp \ + -DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON \ + -DDPCTL_BUILD_CAPI_TESTS=ON \ + -DDPCTL_GENERATE_COVERAGE=ON \ .. make V=1 -n -j 4 && make check && make install diff --git a/dpctl-capi/include/dpctl_sycl_program_interface.h b/dpctl-capi/include/dpctl_sycl_program_interface.h index c931b4b166..8e25d1df40 100644 --- a/dpctl-capi/include/dpctl_sycl_program_interface.h +++ b/dpctl-capi/include/dpctl_sycl_program_interface.h @@ -52,14 +52,17 @@ DPCTL_C_EXTERN_C_BEGIN * @param Ctx An opaque pointer to a sycl::context * @param IL SPIR-V binary * @param Length The size of the IL binary in bytes. + * @param CompileOpts Optional compiler flags used when compiling th + * SPIR-V binary. * @return A new SyclProgramRef pointer if the program creation succeeded, * else returns NULL. */ DPCTL_API __dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx, - __dpctl_keep const void *IL, - size_t Length); +DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef Ctx, + __dpctl_keep const void *IL, + size_t Length, + const char *CompileOpts); /*! * @brief Create a Sycl program from an OpenCL kernel source string. diff --git a/dpctl-capi/source/dpctl_sycl_program_interface.cpp b/dpctl-capi/source/dpctl_sycl_program_interface.cpp index 7f8b75cf63..ae735abbe2 100644 --- a/dpctl-capi/source/dpctl_sycl_program_interface.cpp +++ b/dpctl-capi/source/dpctl_sycl_program_interface.cpp @@ -25,10 +25,15 @@ //===----------------------------------------------------------------------===// #include "dpctl_sycl_program_interface.h" +#include "Config/dpctl_config.h" #include "Support/CBindingWrapping.h" -#include /* Sycl headers */ -#include /* OpenCL headers */ +#include /* Sycl headers */ +#include /* OpenCL headers */ +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION +#include /* Level Zero headers */ +#include +#endif using namespace cl::sycl; @@ -41,7 +46,8 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) __dpctl_give DPCTLSyclProgramRef createOpenCLInterOpProgram (const context &SyclCtx, __dpctl_keep const void *IL, - size_t length) + size_t length, + const char * /* */) { cl_int err; auto CLCtx = SyclCtx.get(); @@ -83,12 +89,63 @@ createOpenCLInterOpProgram (const context &SyclCtx, } } +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION +__dpctl_give DPCTLSyclProgramRef +createLevelZeroInterOpProgram (const context &SyclCtx, + const void *IL, + size_t length, + const char *CompileOpts) +{ + auto ZeCtx = SyclCtx.get_native(); + auto SyclDevices = SyclCtx.get_devices(); + if(SyclDevices.size() > 1) { + // We only support build to one device with Level Zero now. + // TODO: log error + return nullptr; + } + + // Specialization constants are not yet supported. + // Refer https://bit.ly/33UEDYN for details on specialization constants. + ze_module_constants_t ZeSpecConstants = {}; + ZeSpecConstants.numConstants = 0; + + // Populate the Level Zero module descriptions + ze_module_desc_t ZeModuleDesc = {}; + ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; + ZeModuleDesc.inputSize = length; + ZeModuleDesc.pInputModule = (uint8_t*)IL; + ZeModuleDesc.pBuildFlags = CompileOpts; + ZeModuleDesc.pConstants = &ZeSpecConstants; + + auto ZeDevice = SyclDevices[0].get_native(); + ze_module_handle_t ZeModule; + auto ret = zeModuleCreate(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule, + nullptr); + if(ret != ZE_RESULT_SUCCESS) { + // TODO: handle error + return nullptr; + } + + // Create the Sycl program from the ZeModule + try { + auto ZeProgram = new program(sycl::level_zero::make_program( + SyclCtx, reinterpret_cast(ZeModule)) + ); + return wrap(ZeProgram); + } catch (invalid_object_error &e) { + // \todo record error + std::cerr << e.what() << '\n'; + return nullptr; + } +} +#endif } /* end of anonymous namespace */ __dpctl_give DPCTLSyclProgramRef -DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, - __dpctl_keep const void *IL, - size_t length) +DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, + __dpctl_keep const void *IL, + size_t length, + const char *CompileOpts) { DPCTLSyclProgramRef Pref = nullptr; context *SyclCtx = nullptr; @@ -96,21 +153,23 @@ DPCTLProgram_CreateFromOCLSpirv (__dpctl_keep const DPCTLSyclContextRef 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); + Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length, CompileOpts); break; case backend::level_zero: +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + Pref = createLevelZeroInterOpProgram(*SyclCtx, IL, length, + CompileOpts); +#endif break; default: break; } - return Pref; } diff --git a/dpctl-capi/tests/test_sycl_program_interface.cpp b/dpctl-capi/tests/test_sycl_program_interface.cpp index 0e0061b674..791907c1c7 100644 --- a/dpctl-capi/tests/test_sycl_program_interface.cpp +++ b/dpctl-capi/tests/test_sycl_program_interface.cpp @@ -1,4 +1,4 @@ -//===---------- test_sycl_program_interface.cpp - dpctl-C_API --*-- C++ -*-===// +//===---------- test_sycl_program_interface.cpp - dpctl-C_API ---*- C++ -*-===// // // Data Parallel Control Library (dpCtl) // @@ -29,7 +29,7 @@ #include "dpctl_sycl_program_interface.h" #include "dpctl_sycl_queue_interface.h" #include "dpctl_sycl_queue_manager.h" - +#include "Config/dpctl_config.h" #include #include #include @@ -127,12 +127,16 @@ struct TestDPCTLSyclProgramInterface : public ::testing::Test size_t spirvFileSize = 0; std::vector spirvBuffer; size_t nOpenCLGpuQ = 0; +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + size_t nL0GpuQ = 0; +#endif TestDPCTLSyclProgramInterface () : spirvFile{"./multi_kernel.spv", std::ios::binary | std::ios::ate}, spirvFileSize(std::filesystem::file_size("./multi_kernel.spv")), spirvBuffer(spirvFileSize), - nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)) + nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)), + nL0GpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_LEVEL_ZERO, DPCTL_GPU)) { spirvFile.seekg(0, std::ios::beg); spirvFile.read(spirvBuffer.data(), spirvFileSize); @@ -152,7 +156,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource) auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); auto PRef = DPCTLProgram_CreateFromOCLSource(CtxRef, CLProgramStr, - CompileOpts); + CompileOpts); ASSERT_TRUE(PRef != nullptr); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); @@ -162,15 +166,36 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSource) DPCTLProgram_Delete(PRef); } -TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv) +TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvOCL) { if(!nOpenCLGpuQ) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(), - spirvFileSize); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, + nullptr); + ASSERT_TRUE(PRef != nullptr); + ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); + ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); + + DPCTLQueue_Delete(QueueRef); + DPCTLContext_Delete(CtxRef); + DPCTLProgram_Delete(PRef); +} + +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION +TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvL0) +{ + if(!nL0GpuQ) + GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); + + auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_LEVEL_ZERO, DPCTL_GPU, 0); + auto CtxRef = DPCTLQueue_GetContext(QueueRef); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, + nullptr); ASSERT_TRUE(PRef != nullptr); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "add")); ASSERT_TRUE(DPCTLProgram_HasKernel(PRef, "axpy")); @@ -179,6 +204,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromOCLSpirv) DPCTLContext_Delete(CtxRef); DPCTLProgram_Delete(PRef); } +#endif TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource) { @@ -207,15 +233,15 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSource) DPCTLProgram_Delete(PRef); } -TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelOCLSpirv) +TEST_F (TestDPCTLSyclProgramInterface, CheckGetKernelSpirv) { if(!nOpenCLGpuQ) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); auto QueueRef = DPCTLQueueMgr_GetQueue(DPCTL_OPENCL, DPCTL_GPU, 0); auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto PRef = DPCTLProgram_CreateFromOCLSpirv(CtxRef, spirvBuffer.data(), - spirvFileSize); + auto PRef = DPCTLProgram_CreateFromSpirv(CtxRef, spirvBuffer.data(), + spirvFileSize, nullptr); auto AddKernel = DPCTLProgram_GetKernel(PRef, "add"); auto AxpyKernel = DPCTLProgram_GetKernel(PRef, "axpy"); diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index ff6d0c536f..9e9dbca948 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -140,14 +140,15 @@ cdef extern from "dpctl_sycl_context_interface.h": cdef extern from "dpctl_sycl_program_interface.h": - cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSpirv ( + cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromSpirv ( const DPCTLSyclContextRef Ctx, const void *IL, - size_t Length) + size_t Length, + const char *CompileOpts) cdef DPCTLSyclProgramRef DPCTLProgram_CreateFromOCLSource ( const DPCTLSyclContextRef Ctx, - const char* Source, - const char* CompileOpts) + const char *Source, + const char *CompileOpts) cdef DPCTLSyclKernelRef DPCTLProgram_GetKernel (DPCTLSyclProgramRef PRef, const char *KernelName) cdef bool DPCTLProgram_HasKernel (DPCTLSyclProgramRef PRef, diff --git a/dpctl/_sycl_core.pyx b/dpctl/_sycl_core.pyx index 25d01330b2..835df44db1 100644 --- a/dpctl/_sycl_core.pyx +++ b/dpctl/_sycl_core.pyx @@ -678,7 +678,7 @@ cdef class _SyclRTManager: raise UnsupportedDeviceError("Device can only be gpu or cpu") except KeyError: raise UnsupportedBackendError("Backend can only be opencl or " - "level-0") + "level0") def _remove_current_queue(self): DPCTLQueueMgr_PopQueue() @@ -970,7 +970,7 @@ cdef class _SyclRTManager: raise UnsupportedDeviceError("Device can only be gpu or cpu") except KeyError: raise UnsupportedBackendError("Backend can only be opencl or " - "level-0") + "level0") # This private instance of the _SyclQueueManager should not be directly diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd index 27f0716791..1190d5a12c 100644 --- a/dpctl/program/_program.pxd +++ b/dpctl/program/_program.pxd @@ -60,4 +60,5 @@ cdef class SyclProgram: cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*) -cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL) +cpdef create_program_from_spirv (SyclQueue q, const unsigned char[:] IL, + unicode copts=*) diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index efa47f7649..2ea8ef6ea8 100644 --- a/dpctl/program/_program.pyx +++ b/dpctl/program/_program.pyx @@ -162,7 +162,8 @@ cpdef create_program_from_source(SyclQueue q, unicode src, unicode copts=""): cimport cython.array -cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): +cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL, + unicode copts=""): """ Creates a Sycl interoperability program from an SPIR-V binary. @@ -173,6 +174,8 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): q (SyclQueue): The :class:`SyclQueue` for which the :class:`SyclProgram` is going to be built. IL (const char[:]) : SPIR-V binary IL file for an OpenCL program. + copts (unicode) : Optional compilation flags that will be used + when compiling the program. Returns: program (SyclProgram): A :class:`SyclProgram` object wrapping the sycl::program returned by the C API. @@ -185,7 +188,9 @@ cpdef create_program_from_spirv(SyclQueue q, const unsigned char[:] IL): cdef const unsigned char *dIL = &IL[0] cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref() cdef size_t length = IL.shape[0] - Pref = DPCTLProgram_CreateFromOCLSpirv(CRef, dIL, length) + cdef bytes bCOpts = copts.encode('utf8') + cdef const char *COpts = bCOpts + Pref = DPCTLProgram_CreateFromSpirv(CRef, dIL, length, COpts) if Pref is NULL: raise SyclProgramCompilationError() diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index 5eb1c2b66b..ce7f6051f2 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -87,9 +87,16 @@ def test_create_program_from_spirv(self): "No Level0 GPU queues available", ) class TestProgramForLevel0GPU(unittest.TestCase): - @unittest.expectedFailure - def test_create_program_from_spirv(self): + import sys + + # Level zero program creation from a SPIR-V binary is not supported + # on Windows. + @unittest.skipIf( + sys.platform in ["win32", "cygwin"], + "Level Zero module creation unsupported on Windows.", + ) + 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: diff --git a/scripts/build_backend.py b/scripts/build_backend.py index 2708094dd6..c68dd37b40 100644 --- a/scripts/build_backend.py +++ b/scripts/build_backend.py @@ -43,6 +43,7 @@ "-DDPCPP_INSTALL_DIR=" + DPCPP_ROOT, "-DCMAKE_C_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang"), "-DCMAKE_CXX_COMPILER:PATH=" + os.path.join(DPCPP_ROOT, "bin", "clang++"), + "-DDPCTL_ENABLE_LO_PROGRAM_CREATION=ON", backends, ] subprocess.check_call(cmake_args, stderr=subprocess.STDOUT, shell=False) From 57d41faebdec22692e87cfafafe44832fb85ea9d Mon Sep 17 00:00:00 2001 From: Diptorup Deb <3046810+diptorupd@users.noreply.github.com> Date: Sat, 23 Jan 2021 23:44:44 -0600 Subject: [PATCH 5/8] Create CODE_OF_CONDUCT.md (#251) --- CODE_OF_CONDUCT.md | 76 ++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 76 insertions(+) create mode 100644 CODE_OF_CONDUCT.md diff --git a/CODE_OF_CONDUCT.md b/CODE_OF_CONDUCT.md new file mode 100644 index 0000000000..2ccf066297 --- /dev/null +++ b/CODE_OF_CONDUCT.md @@ -0,0 +1,76 @@ +# Contributor Covenant Code of Conduct + +## Our Pledge + +In the interest of fostering an open and welcoming environment, we as +contributors and maintainers pledge to making participation in our project and +our community a harassment-free experience for everyone, regardless of age, body +size, disability, ethnicity, sex characteristics, gender identity and expression, +level of experience, education, socio-economic status, nationality, personal +appearance, race, religion, or sexual identity and orientation. + +## Our Standards + +Examples of behavior that contributes to creating a positive environment +include: + +* Using welcoming and inclusive language +* Being respectful of differing viewpoints and experiences +* Gracefully accepting constructive criticism +* Focusing on what is best for the community +* Showing empathy towards other community members + +Examples of unacceptable behavior by participants include: + +* The use of sexualized language or imagery and unwelcome sexual attention or + advances +* Trolling, insulting/derogatory comments, and personal or political attacks +* Public or private harassment +* Publishing others' private information, such as a physical or electronic + address, without explicit permission +* Other conduct which could reasonably be considered inappropriate in a + professional setting + +## Our Responsibilities + +Project maintainers are responsible for clarifying the standards of acceptable +behavior and are expected to take appropriate and fair corrective action in +response to any instances of unacceptable behavior. + +Project maintainers have the right and responsibility to remove, edit, or +reject comments, commits, code, wiki edits, issues, and other contributions +that are not aligned to this Code of Conduct, or to ban temporarily or +permanently any contributor for other behaviors that they deem inappropriate, +threatening, offensive, or harmful. + +## Scope + +This Code of Conduct applies both within project spaces and in public spaces +when an individual is representing the project or its community. Examples of +representing a project or community include using an official project e-mail +address, posting via an official social media account, or acting as an appointed +representative at an online or offline event. Representation of a project may be +further defined and clarified by project maintainers. + +## Enforcement + +Instances of abusive, harassing, or otherwise unacceptable behavior may be +reported by contacting the project team at scripting@intel.com. All +complaints will be reviewed and investigated and will result in a response that +is deemed necessary and appropriate to the circumstances. The project team is +obligated to maintain confidentiality with regard to the reporter of an incident. +Further details of specific enforcement policies may be posted separately. + +Project maintainers who do not follow or enforce the Code of Conduct in good +faith may face temporary or permanent repercussions as determined by other +members of the project's leadership. + +## Attribution + +This Code of Conduct is adapted from the [Contributor Covenant][homepage], version 1.4, +available at https://www.contributor-covenant.org/version/1/4/code-of-conduct.html + +[homepage]: https://www.contributor-covenant.org + +For answers to common questions about this code of conduct, see +https://www.contributor-covenant.org/faq From d10dd68d87e24f01f4266bc5d7f0ce3d8d943782 Mon Sep 17 00:00:00 2001 From: Diptorup Deb <3046810+diptorupd@users.noreply.github.com> Date: Mon, 25 Jan 2021 22:19:10 -0600 Subject: [PATCH 6/8] Level zero codegen (#250) * WIP level zero support. * Add support to build a SYCL program from SPIR-V. * Ignore the generated config file. * Skip test on Windows. * Black formatting. * Skip test properly. * Fix comment. * Fix black error. * Fix comments. * Fix potential issue on Windows. * Make level zero a dynamically loaded library. * Do not install helper headers and run tests during linux build. * Cleanups... --- dpctl-capi/CMakeLists.txt | 9 -- .../helper/include/dpctl_dynamic_lib_helper.h | 108 ++++++++++++++++++ .../source/dpctl_sycl_program_interface.cpp | 56 ++++++++- .../tests/test_sycl_program_interface.cpp | 7 +- 4 files changed, 160 insertions(+), 20 deletions(-) create mode 100644 dpctl-capi/helper/include/dpctl_dynamic_lib_helper.h diff --git a/dpctl-capi/CMakeLists.txt b/dpctl-capi/CMakeLists.txt index 16748f66cf..4fe98c7c39 100644 --- a/dpctl-capi/CMakeLists.txt +++ b/dpctl-capi/CMakeLists.txt @@ -94,9 +94,6 @@ if(DPCTL_ENABLE_LO_PROGRAM_CREATION) PRIVATE ${LEVEL_ZERO_INCLUDE_DIR} ) - target_link_libraries(DPCTLSyclInterface - PRIVATE ${LEVEL_ZERO_LIBRARY} - ) else() message(WARNING "DPCTL support Level Zero program creation not supported " @@ -124,12 +121,6 @@ foreach(HEADER ${HEADERS}) install(FILES "${HEADER}" DESTINATION include/Support) endforeach() -# Install all headers in helper/include -file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/helper/include/*.h") -foreach(HEADER ${HEADERS}) - install(FILES "${HEADER}" DESTINATION helper/include) -endforeach() - # Install all headers in include/Config file(GLOB HEADERS "${CMAKE_SOURCE_DIR}/include/Config/*.h") foreach(HEADER ${HEADERS}) diff --git a/dpctl-capi/helper/include/dpctl_dynamic_lib_helper.h b/dpctl-capi/helper/include/dpctl_dynamic_lib_helper.h new file mode 100644 index 0000000000..34abce2434 --- /dev/null +++ b/dpctl-capi/helper/include/dpctl_dynamic_lib_helper.h @@ -0,0 +1,108 @@ +//===--------------- dpctl_dynamic_lib_helper.h - dpctl-C_API -*-C++-*-===// +// +// Data Parallel Control Library (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 +/// Helper for dynamic libs management. +//===----------------------------------------------------------------------===// + +#ifndef __DPCTL_DYNAMIC_LIB_HELPER_H__ +#define __DPCTL_DYNAMIC_LIB_HELPER_H__ + +#if defined(__linux__) || defined(_WIN32) || defined(_WIN64) + + #ifdef __linux__ + + #include + + #elif defined(_WIN32) || defined(_WIN64) + + #define NOMINMAX + #include + + #endif // __linux__ + +#include + +namespace dpctl +{ + +class DynamicLibHelper final +{ +public: + DynamicLibHelper() = delete; + DynamicLibHelper(const DynamicLibHelper &) = delete; + DynamicLibHelper(const char * libName, int flag) + { + + #ifdef __linux__ + _handle = dlopen(libName, flag); + #elif defined(_WIN32) || defined(_WIN64) + _handle = LoadLibraryA(libName); + #endif + } + + ~DynamicLibHelper() + { + #ifdef __linux__ + dlclose(_handle); + #elif defined(_WIN32) || defined(_WIN64) + FreeLibrary((HMODULE)_handle); + #endif + }; + + template + T getSymbol(const char * symName) + { + #ifdef __linux__ + void * sym = dlsym(_handle, symName); + char * error = dlerror(); + + if (NULL != error) + { + return nullptr; + } + #elif defined(_WIN32) || defined(_WIN64) + void * sym = (void *)GetProcAddress((HMODULE)_handle, symName); + + if (NULL == sym) + { + return nullptr; + } + #endif + + return (T)sym; + } + + bool opened () const + { + if (!_handle) + return false; + else + return true; + } + +private: + void * _handle = nullptr; +}; + +} // namespace dpctl + +#endif // #if defined(__linux__) || defined(_WIN32) || defined(_WIN64) +#endif // __DPCTL_DYNAMIC_LIB_HELPER_H__ diff --git a/dpctl-capi/source/dpctl_sycl_program_interface.cpp b/dpctl-capi/source/dpctl_sycl_program_interface.cpp index ae735abbe2..d4d0477acb 100644 --- a/dpctl-capi/source/dpctl_sycl_program_interface.cpp +++ b/dpctl-capi/source/dpctl_sycl_program_interface.cpp @@ -33,12 +33,32 @@ #ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION #include /* Level Zero headers */ #include +#include "../helper/include/dpctl_dynamic_lib_helper.h" #endif using namespace cl::sycl; namespace { +#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + +#ifdef __linux__ +static const char * zeLoaderName = "libze_loader.so"; +static const int libLoadFlags = RTLD_NOLOAD | RTLD_NOW | RTLD_LOCAL; +#else + #error "Level Zero program compilation is unavailable for this platform" +#endif + +typedef ze_result_t (*zeModuleCreateFT)(ze_context_handle_t, + ze_device_handle_t, + const ze_module_desc_t *, + ze_module_handle_t *, + ze_module_build_log_handle_t *); + +const char * zeModuleCreateFuncName = "zeModuleCreate"; + +#endif // #ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + DEFINE_SIMPLE_CONVERSION_FUNCTIONS(context, DPCTLSyclContextRef) DEFINE_SIMPLE_CONVERSION_FUNCTIONS(program, DPCTLSyclProgramRef) DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) @@ -90,6 +110,23 @@ createOpenCLInterOpProgram (const context &SyclCtx, } #ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION + +zeModuleCreateFT getZeModuleCreateFn () +{ + static dpctl::DynamicLibHelper zeLib(zeLoaderName, libLoadFlags); + if(!zeLib.opened()) { + // TODO: handle error + std::cerr << "The level zero loader dynamic library could not " + "be opened.\n"; + return nullptr; + } + static auto stZeModuleCreateF = zeLib.getSymbol( + zeModuleCreateFuncName + ); + + return stZeModuleCreateF; +} + __dpctl_give DPCTLSyclProgramRef createLevelZeroInterOpProgram (const context &SyclCtx, const void *IL, @@ -99,8 +136,8 @@ createLevelZeroInterOpProgram (const context &SyclCtx, auto ZeCtx = SyclCtx.get_native(); auto SyclDevices = SyclCtx.get_devices(); if(SyclDevices.size() > 1) { - // We only support build to one device with Level Zero now. - // TODO: log error + std::cerr << "Level zero program can be created for only one device.\n"; + // TODO: handle error return nullptr; } @@ -119,8 +156,14 @@ createLevelZeroInterOpProgram (const context &SyclCtx, auto ZeDevice = SyclDevices[0].get_native(); ze_module_handle_t ZeModule; - auto ret = zeModuleCreate(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule, - nullptr); + + auto stZeModuleCreateF = getZeModuleCreateFn(); + + if(!stZeModuleCreateF) + return nullptr; + + auto ret = stZeModuleCreateF(ZeCtx, ZeDevice, &ZeModuleDesc, &ZeModule, + nullptr); if(ret != ZE_RESULT_SUCCESS) { // TODO: handle error return nullptr; @@ -138,7 +181,8 @@ createLevelZeroInterOpProgram (const context &SyclCtx, return nullptr; } } -#endif +#endif /* #ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION */ + } /* end of anonymous namespace */ __dpctl_give DPCTLSyclProgramRef @@ -156,7 +200,7 @@ DPCTLProgram_CreateFromSpirv (__dpctl_keep const DPCTLSyclContextRef CtxRef, SyclCtx = unwrap(CtxRef); // get the backend type auto BE = SyclCtx->get_platform().get_backend(); - switch (BE) + switch(BE) { case backend::opencl: Pref = createOpenCLInterOpProgram(*SyclCtx, IL, length, CompileOpts); diff --git a/dpctl-capi/tests/test_sycl_program_interface.cpp b/dpctl-capi/tests/test_sycl_program_interface.cpp index 791907c1c7..7a7f49819f 100644 --- a/dpctl-capi/tests/test_sycl_program_interface.cpp +++ b/dpctl-capi/tests/test_sycl_program_interface.cpp @@ -127,16 +127,12 @@ struct TestDPCTLSyclProgramInterface : public ::testing::Test size_t spirvFileSize = 0; std::vector spirvBuffer; size_t nOpenCLGpuQ = 0; -#ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION - size_t nL0GpuQ = 0; -#endif TestDPCTLSyclProgramInterface () : spirvFile{"./multi_kernel.spv", std::ios::binary | std::ios::ate}, spirvFileSize(std::filesystem::file_size("./multi_kernel.spv")), spirvBuffer(spirvFileSize), - nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)), - nL0GpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_LEVEL_ZERO, DPCTL_GPU)) + nOpenCLGpuQ(DPCTLQueueMgr_GetNumQueues(DPCTL_OPENCL, DPCTL_GPU)) { spirvFile.seekg(0, std::ios::beg); spirvFile.read(spirvBuffer.data(), spirvFileSize); @@ -188,6 +184,7 @@ TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvOCL) #ifdef DPCTL_ENABLE_LO_PROGRAM_CREATION TEST_F (TestDPCTLSyclProgramInterface, CheckCreateFromSpirvL0) { + auto nL0GpuQ = DPCTLQueueMgr_GetNumQueues(DPCTL_LEVEL_ZERO, DPCTL_GPU); if(!nL0GpuQ) GTEST_SKIP_("Skipping as no OpenCL GPU device found.\n"); From e3c0db00a4195a3143dc7567505a79b88b052b79 Mon Sep 17 00:00:00 2001 From: Sergey Pokhodenko Date: Fri, 29 Jan 2021 06:46:01 -0600 Subject: [PATCH 7/8] Update CHANGELOG --- CHANGELOG.md | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index beb1b60507..2db11d48f5 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,6 +7,16 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ## [Unreleased] +### Added +- Documentation improvements +- Cmake improvements and Coverage for C API +- Add support for Level Zero + +### Fixed +- Remove `cython` from `install_requires`. It allows use `dpCtl` in `numba` extensions. +- Incorrect import in example. + + ## [0.5.0] - 2020-12-17 ### Added - `_Memory.get_pointer_type` static method which returns kind of USM pointer. From be7b0b9317a0998adcb20e1f04727c8975408fd3 Mon Sep 17 00:00:00 2001 From: Sergey Pokhodenko Date: Fri, 29 Jan 2021 06:47:24 -0600 Subject: [PATCH 8/8] Update CHANGELOG --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 2db11d48f5..5479c3109d 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,6 +11,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 - Documentation improvements - Cmake improvements and Coverage for C API - Add support for Level Zero +- Code of conduct ### Fixed - Remove `cython` from `install_requires`. It allows use `dpCtl` in `numba` extensions.