diff --git a/CMakeLists.txt b/CMakeLists.txt index 6f80d05dc8a..5c339fb0abd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -5,7 +5,7 @@ if (CMAKE_VERSION VERSION_GREATER_EQUAL 3.12) cmake_policy(SET CMP0074 NEW) endif() -# Let CAS handle the CUDA architecture flags (for now) +# Let CAS handle the CUDA architecture flags(for now) # Windows still gives CMP0104 warning if putting it in cuda. if (CMAKE_VERSION VERSION_GREATER_EQUAL 3.18) cmake_policy(SET CMP0104 OLD) @@ -27,6 +27,7 @@ option(GINKGO_BUILD_EXAMPLES "Build Ginkgo's examples" ON) option(GINKGO_BUILD_BENCHMARKS "Build Ginkgo's benchmarks" ON) option(GINKGO_BUILD_REFERENCE "Compile reference CPU kernels" ON) option(GINKGO_BUILD_OMP "Compile OpenMP kernels for CPU" ${GINKGO_HAS_OMP}) +option(GINKGO_BUILD_DPCPP "Compile DPC++ kernels for CPU" ${GINKGO_HAS_DPCPP}) option(GINKGO_BUILD_CUDA "Compile kernels for NVIDIA GPUs" ${GINKGO_HAS_CUDA}) option(GINKGO_BUILD_HIP "Compile kernels for AMD or NVIDIA GPUs" ${GINKGO_HAS_HIP}) option(GINKGO_BUILD_DOC "Generate documentation" OFF) @@ -72,7 +73,8 @@ option(BUILD_SHARED_LIBS "Build shared (.so, .dylib, .dll) libraries" ON) set(GINKGO_CIRCULAR_DEPS_FLAGS "-Wl,--no-undefined") if(BUILD_SHARED_LIBS AND (WIN32 OR CYGWIN) AND (GINKGO_BUILD_TESTS OR GINKGO_BUILD_EXAMPLES OR GINKGO_BUILD_BENCHMARKS)) - # Change shared libraries output only if this build has executable program with shared libraries. + # Change shared libraries output only if this build has executable program + # with shared libraries. set(GINKGO_CHANGED_SHARED_LIBRARY TRUE) option(GINKGO_CHECK_PATH "Tell Ginkgo to check if the environment variable PATH is available for this build." ON) set(GINKGO_WINDOWS_SHARED_LIBRARY_RELPATH "windows_shared_library" CACHE STRING @@ -83,7 +85,7 @@ else() set(GINKGO_CHANGED_SHARED_LIBRARY FALSE) endif() -if(GINKGO_BUILD_TESTS AND (GINKGO_BUILD_CUDA OR GINKGO_BUILD_OMP OR GINKGO_BUILD_HIP)) +if(GINKGO_BUILD_TESTS AND (GINKGO_BUILD_CUDA OR GINKGO_BUILD_OMP OR GINKGO_BUILD_HIP OR GINKGO_BUILD_DPCPP)) message(STATUS "GINKGO_BUILD_TESTS is ON, enabling GINKGO_BUILD_REFERENCE") set(GINKGO_BUILD_REFERENCE ON CACHE BOOL "Compile reference CPU kernels" FORCE) endif() @@ -105,12 +107,12 @@ if(NOT DEFINED CMAKE_DEBUG_POSTFIX) endif() if(GINKGO_BUILD_TESTS) - # Configure CTest +# Configure CTest configure_file( ${CMAKE_CURRENT_LIST_DIR}/cmake/CTestCustom.cmake.in ${CMAKE_CURRENT_BINARY_DIR}/CTestCustom.cmake @ONLY) - # For testing, we need some special matrices +#For testing, we need some special matrices add_subdirectory(matrices) enable_testing() @@ -127,8 +129,9 @@ endif() list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake/Modules/") -# Find important header files, store the definitions in include/ginkgo/config.h.in -# For details, see https://gitlab.kitware.com/cmake/community/wikis/doc/tutorials/How-To-Write-Platform-Checks +# Find important header files, store the definitions in +# include/ginkgo/config.h.in For details, see +# https://gitlab.kitware.com/cmake/community/wikis/doc/tutorials/How-To-Write-Platform-Checks include(CheckIncludeFileCXX) check_include_file_cxx(cxxabi.h GKO_HAVE_CXXABI_H) @@ -143,11 +146,12 @@ set(GINKGO_HIP_PLATFORM_NVCC 0) set(GINKGO_HIP_PLATFORM_HCC 0) if(GINKGO_BUILD_HIP) - # GINKGO_HIPCONFIG_PATH and HIP_PATH are set in cmake/hip_path.cmake +# GINKGO_HIPCONFIG_PATH and HIP_PATH are set in cmake/hip_path.cmake if(DEFINED ENV{HIP_PLATFORM}) set(GINKGO_HIP_PLATFORM "$ENV{HIP_PLATFORM}") elseif(GINKGO_HIPCONFIG_PATH) - execute_process(COMMAND ${GINKGO_HIPCONFIG_PATH} --platform OUTPUT_VARIABLE GINKGO_HIP_PLATFORM) + execute_process(COMMAND ${ + GINKGO_HIPCONFIG_PATH} --platform OUTPUT_VARIABLE GINKGO_HIP_PLATFORM) else() message(FATAL_ERROR "No platform could be found for HIP. " "Set and export the environment variable HIP_PLATFORM.") @@ -170,7 +174,8 @@ include(cmake/hip_helpers.cmake) include(cmake/install_helpers.cmake) include(cmake/windows_helpers.cmake) -# This is modified from https://gitlab.kitware.com/cmake/community/wikis/FAQ#dynamic-replace +# This is modified from +# https://gitlab.kitware.com/cmake/community/wikis/FAQ#dynamic-replace if(MSVC) if(BUILD_SHARED_LIBS) ginkgo_switch_to_windows_dynamic("CXX") @@ -204,12 +209,16 @@ endif() if (GINKGO_BUILD_OMP) add_subdirectory(omp) # High-performance omp kernels endif() +if (GINKGO_BUILD_DPCPP) + add_subdirectory(dpcpp) # High-performance DPC++ kernels +endif() # HIP needs to be last because it builds the GINKGO_RPATH_FOR_HIP variable # which needs to know the `ginkgo` target. if(GINKGO_BUILD_HIP) add_subdirectory(hip) # High-performance kernels for AMD or NVIDIA GPUs endif() + # Non core directories and targets if(GINKGO_BUILD_EXAMPLES) add_subdirectory(examples) @@ -226,8 +235,8 @@ if(GINKGO_DEVEL_TOOLS) add_dependencies(format add_license) endif() -# MacOS needs to install bash, gnu-sed, findutils and coreutils -# format_header needs clang-format 6.0.0+ +# MacOS needs to install bash, gnu - sed, findutils and coreutils +# format_header needs clang - format 6.0.0 + find_program(BASH bash) if(NOT "${BASH}" STREQUAL "BASH-NOTFOUND" AND GINKGO_DEVEL_TOOLS) add_custom_target(generate_ginkgo_header ALL @@ -271,14 +280,15 @@ endif() configure_file(${Ginkgo_SOURCE_DIR}/cmake/ginkgo.pc.in ${Ginkgo_BINARY_DIR}/ginkgo.pc @ONLY) -# WINDOWS NVCC has " inside the string, add escape charater to avoid config problem. +# WINDOWS NVCC has " inside the string, add escape character +# to avoid config problem. ginkgo_modify_flags(CMAKE_CUDA_FLAGS) ginkgo_modify_flags(CMAKE_CUDA_FLAGS_DEBUG) ginkgo_modify_flags(CMAKE_CUDA_FLAGS_RELEASE) ginkgo_install() if(MSVC) - # Set path/command with $ +# Set path / command with $ < CONFIG> set(GINKGO_TEST_INSTALL_COMMAND "${Ginkgo_BINARY_DIR}/test_install/$/test_install") if(GINKGO_BUILD_CUDA) set(GINKGO_TEST_INSTALL_COMMAND "${GINKGO_TEST_INSTALL_COMMAND}" "${Ginkgo_BINARY_DIR}/test_install/$/test_install_cuda") @@ -295,8 +305,9 @@ add_custom_target(test_install -DCMAKE_PREFIX_PATH=${CMAKE_INSTALL_PREFIX}/${GINKGO_INSTALL_CONFIG_DIR} -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER} -DCMAKE_CUDA_COMPILER=${CMAKE_CUDA_COMPILER} - # `--config cfg` is ignored by single-configuration generator. - # `$` is always be the same as `CMAKE_BUILD_TYPE` in single-configuration generator. +# `-- config cfg` is ignored by single - configuration generator. +# `$ ` is always be the same as `CMAKE_BUILD_TYPE` in single - \ + configuration generator. COMMAND ${CMAKE_COMMAND} --build ${Ginkgo_BINARY_DIR}/test_install --config $ COMMAND ${GINKGO_TEST_INSTALL_COMMAND} COMMENT "Running a test on the installed binaries. This requires running `(sudo) make install` first.") @@ -309,7 +320,6 @@ set(CPACK_PACKAGE_CONTACT "ginkgo.library@gmail.com") include(CPack) # And finally, print the configuration to screen: -# if(GINKGO_CONFIG_LOG_DETAILED) FILE(READ ${PROJECT_BINARY_DIR}/detailed.log GINKGO_LOG_SUMMARY) else() diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index d3ec0799488..899311fd6f3 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -68,8 +68,8 @@ look at our coding guidelines before proposing a pull request. Ginkgo is divided into a `core` module with common functionalities independent of the architecture, and several kernel modules (`reference`, `omp`, `cuda`, -`hip`) which contain low-level computational routines for each supported -architecture. +`hip`, `dpcpp`) which contain low-level computational routines for each +supported architecture. ### Extended header files diff --git a/INSTALL.md b/INSTALL.md index 2a0056bf3a4..7ef20b65aef 100644 --- a/INSTALL.md +++ b/INSTALL.md @@ -32,6 +32,10 @@ Ginkgo adds the following additional switches to control what is being built: * `-DGINKGO_BUILD_CUDA={ON, OFF}` builds optimized cuda versions of the kernels (requires CUDA), default is `ON` if a CUDA compiler could be detected, `OFF` otherwise. +* `-DGINKGO_BUILD_DPCPP={ON, OFF}` builds optimized DPC++ versions of the + kernels (requires `CMAKE_CXX_COMPILER` to be set to the `dpcpp` compiler). + The default is `ON` if `CMAKE_CXX_COMPILER` is a DPC++ compiler, `OFF` + otherwise. * `-DGINKGO_BUILD_HIP={ON, OFF}` builds optimized HIP versions of the kernels (requires HIP), default is `ON` if an installation of HIP could be detected, `OFF` otherwise. diff --git a/benchmark/utils/overhead_linop.hpp b/benchmark/utils/overhead_linop.hpp index 2c2aa4baccf..5bddad88325 100644 --- a/benchmark/utils/overhead_linop.hpp +++ b/benchmark/utils/overhead_linop.hpp @@ -109,6 +109,15 @@ GKO_DECLARE_ALL; } // namespace hip +namespace dpcpp { +namespace overhead { + +GKO_DECLARE_ALL; + +} // namespace overhead +} // namespace dpcpp + + #undef GKO_DECLARE_ALL diff --git a/cmake/GinkgoConfig.cmake.in b/cmake/GinkgoConfig.cmake.in index 0348f956e7b..da1776a3d1b 100644 --- a/cmake/GinkgoConfig.cmake.in +++ b/cmake/GinkgoConfig.cmake.in @@ -36,6 +36,7 @@ set(GINKGO_BUILD_REFERENCE @GINKGO_BUILD_REFERENCE@) set(GINKGO_BUILD_OMP @GINKGO_BUILD_OMP@) set(GINKGO_BUILD_CUDA @GINKGO_BUILD_CUDA@) set(GINKGO_BUILD_HIP @GINKGO_BUILD_HIP@) +set(GINKGO_BUILD_DPCPP @GINKGO_BUILD_HIP@) set(GINKGO_DEVEL_TOOLS @GINKGO_DEVEL_TOOLS@) set(GINKGO_BUILD_TESTS @GINKGO_BUILD_TESTS@) diff --git a/cmake/autodetect_executors.cmake b/cmake/autodetect_executors.cmake index 1f90640acb9..c432a92d593 100644 --- a/cmake/autodetect_executors.cmake +++ b/cmake/autodetect_executors.cmake @@ -1,9 +1,13 @@ set(GINKGO_HAS_OMP OFF) set(GINKGO_HAS_CUDA OFF) +set(GINKGO_HAS_DPCPP OFF) set(GINKGO_HAS_HIP OFF) find_package(OpenMP) include(CheckLanguage) check_language(CUDA) +try_compile(GKO_CAN_COMPILE_DPCPP ${PROJECT_BINARY_DIR}/dpcpp + SOURCES ${PROJECT_SOURCE_DIR}/dpcpp/test_dpcpp.dp.cpp + CXX_STANDARD 17) if(OpenMP_CXX_FOUND) if(NOT DEFINED GINKGO_BUILD_OMP) @@ -25,3 +29,10 @@ if(GINKGO_HIPCONFIG_PATH) endif() set(GINKGO_HAS_HIP ON) endif() + +if (GKO_CAN_COMPILE_DPCPP) + if(NOT DEFINED GINKGO_BUILD_DPCPP) + message(STATUS "Enabling DPCPP executor") + endif() + set(GINKGO_HAS_DPCPP ON) +endif() diff --git a/cmake/create_test.cmake b/cmake/create_test.cmake index f2ff8496e64..a2d2ef5d978 100644 --- a/cmake/create_test.cmake +++ b/cmake/create_test.cmake @@ -17,6 +17,26 @@ function(ginkgo_create_test test_name) add_test(NAME ${REL_BINARY_DIR}/${test_name} COMMAND ${TEST_TARGET_NAME}) endfunction(ginkgo_create_test) +function(ginkgo_create_dpcpp_test test_name) + file(RELATIVE_PATH REL_BINARY_DIR + ${PROJECT_BINARY_DIR} ${CMAKE_CURRENT_BINARY_DIR}) + string(REPLACE "/" "_" TEST_TARGET_NAME "${REL_BINARY_DIR}/${test_name}") + add_executable(${TEST_TARGET_NAME} ${test_name}.dp.cpp) + target_compile_features("${TEST_TARGET_NAME}" PUBLIC cxx_std_17) + target_compile_options("${TEST_TARGET_NAME}" PRIVATE "${GINKGO_DPCPP_FLAGS}") + target_include_directories("${TEST_TARGET_NAME}" + PRIVATE + "$" + ) + set_target_properties(${TEST_TARGET_NAME} PROPERTIES + OUTPUT_NAME ${test_name}) + if (GINKGO_CHECK_CIRCULAR_DEPS) + target_link_libraries(${TEST_TARGET_NAME} PRIVATE "${GINKGO_CIRCULAR_DEPS_FLAGS}") + endif() + target_link_libraries(${TEST_TARGET_NAME} PRIVATE ginkgo GTest::Main GTest::GTest ${ARGN}) + add_test(NAME ${REL_BINARY_DIR}/${test_name} COMMAND ${TEST_TARGET_NAME}) +endfunction(ginkgo_create_dpcpp_test) + function(ginkgo_create_thread_test test_name) set(THREADS_PREFER_PTHREAD_FLAG ON) find_package(Threads REQUIRED) diff --git a/cmake/get_info.cmake b/cmake/get_info.cmake index f4773be50bd..94f898a24c4 100644 --- a/cmake/get_info.cmake +++ b/cmake/get_info.cmake @@ -72,6 +72,12 @@ function(ginkgo_print_variable log_type var_name) FILE(APPEND ${log_type} "${upd_string}") endfunction() +macro(ginkgo_print_foreach_variable variables) + foreach(var ${variables}) + ginkgo_print_variable(${${log_type}} ${var} ) + endforeach() +endmacro() + IF("${GINKGO_GIT_SHORTREV}" STREQUAL "") set(to_print "Summary of Configuration for Ginkgo (version ${Ginkgo_VERSION} with tag ${Ginkgo_VERSION_TAG}) --" @@ -88,12 +94,9 @@ ENDIF() set(log_types "detailed_log;minimal_log") foreach(log_type ${log_types}) ginkgo_print_module_footer(${${log_type}} "Ginkgo configuration:") - set(print_var - "CMAKE_BUILD_TYPE;BUILD_SHARED_LIBS;CMAKE_INSTALL_PREFIX;PROJECT_SOURCE_DIR;PROJECT_BINARY_DIR" - ) - foreach(var ${print_var}) - ginkgo_print_variable(${${log_type}} ${var} ) - endforeach() + ginkgo_print_foreach_variable( + "CMAKE_BUILD_TYPE;BUILD_SHARED_LIBS;CMAKE_INSTALL_PREFIX" + "PROJECT_SOURCE_DIR;PROJECT_BINARY_DIR") string(SUBSTRING " -- CMAKE_CXX_COMPILER: " 0 55 print_string) @@ -108,31 +111,17 @@ foreach(log_type ${log_types}) FILE(APPEND ${${log_type}} "${print_string}") ginkgo_print_module_footer(${${log_type}} "User configuration:") ginkgo_print_module_footer(${${log_type}} " Enabled modules:") - set(print_var - "GINKGO_BUILD_OMP;GINKGO_BUILD_REFERENCE;GINKGO_BUILD_CUDA;GINKGO_BUILD_HIP" - ) - foreach(var ${print_var}) - ginkgo_print_variable(${${log_type}} ${var} ) - endforeach() + ginkgo_print_foreach_variable( + "GINKGO_BUILD_OMP;GINKGO_BUILD_REFERENCE;GINKGO_BUILD_CUDA;GINKGO_BUILD_HIP;GINKGO_BUILD_DPCPP") ginkgo_print_module_footer(${${log_type}} " Tests, benchmarks and examples:") - set(print_var + ginkgo_print_foreach_variable( "GINKGO_BUILD_TESTS;GINKGO_BUILD_EXAMPLES;GINKGO_EXTLIB_EXAMPLE;GINKGO_BUILD_BENCHMARKS") - foreach(var ${print_var}) - ginkgo_print_variable(${${log_type}} ${var} ) - endforeach() ginkgo_print_module_footer(${${log_type}} " Documentation:") - set(print_var - "GINKGO_BUILD_DOC;GINKGO_VERBOSE_LEVEL") - foreach(var ${print_var}) - ginkgo_print_variable(${${log_type}} ${var} ) - endforeach() + ginkgo_print_foreach_variable("GINKGO_BUILD_DOC;GINKGO_VERBOSE_LEVEL") ginkgo_print_module_footer(${${log_type}} " Developer helpers:") - set(print_var + ginkgo_print_foreach_variable( "GINKGO_DEVEL_TOOLS;GINKGO_WITH_CLANG_TIDY;GINKGO_WITH_IWYU" "GINKGO_CHECK_CIRCULAR_DEPS;GINKGO_CHECK_PATH") - foreach(var ${print_var}) - ginkgo_print_variable(${${log_type}} ${var} ) - endforeach() ginkgo_print_module_footer(${${log_type}} "") endforeach() @@ -160,6 +149,10 @@ IF(GINKGO_BUILD_HIP) include(hip/get_info.cmake) ENDIF() +IF(GINKGO_BUILD_DPCPP) + include(dpcpp/get_info.cmake) +ENDIF() + ginkgo_print_generic_header(${detailed_log} "Optional Components:") ginkgo_print_variable(${detailed_log} "GKO_HAVE_PAPI_SDE") if(PAPI_sde_FOUND) diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index f5cff2e407d..41c2087a25b 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -57,7 +57,7 @@ target_compile_options(ginkgo PRIVATE "${GINKGO_COMPILER_FLAGS}") # regardless of whether it is installed or added as a subdirectory add_library(Ginkgo::ginkgo ALIAS ginkgo) target_link_libraries(ginkgo - PUBLIC ginkgo_omp ginkgo_cuda ginkgo_reference ginkgo_hip) + PUBLIC ginkgo_omp ginkgo_cuda ginkgo_reference ginkgo_hip ginkgo_dpcpp) # The PAPI dependency needs to be exposed to the user. if (GINKGO_HAVE_PAPI_SDE) target_link_libraries(ginkgo PUBLIC PAPI::PAPI) diff --git a/core/base/executor.cpp b/core/base/executor.cpp index 9d80ad818f0..90342b30e43 100644 --- a/core/base/executor.cpp +++ b/core/base/executor.cpp @@ -53,6 +53,10 @@ void Operation::run(std::shared_ptr executor) const GKO_NOT_IMPLEMENTED; +void Operation::run(std::shared_ptr executor) const + GKO_NOT_IMPLEMENTED; + + void Operation::run(std::shared_ptr executor) const { this->run(static_cast>(executor)); diff --git a/core/components/absolute_array.hpp b/core/components/absolute_array.hpp index d48d04335fa..746c1de733a 100644 --- a/core/components/absolute_array.hpp +++ b/core/components/absolute_array.hpp @@ -99,6 +99,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/components/fill_array.hpp b/core/components/fill_array.hpp index 7bafb8aecb4..452661bed84 100644 --- a/core/components/fill_array.hpp +++ b/core/components/fill_array.hpp @@ -91,6 +91,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/components/precision_conversion.hpp b/core/components/precision_conversion.hpp index 719c596c34e..82b684d4be1 100644 --- a/core/components/precision_conversion.hpp +++ b/core/components/precision_conversion.hpp @@ -93,6 +93,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/components/prefix_sum.hpp b/core/components/prefix_sum.hpp index d171be831aa..f226ce4c282 100644 --- a/core/components/prefix_sum.hpp +++ b/core/components/prefix_sum.hpp @@ -91,6 +91,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/device_hooks/CMakeLists.txt b/core/device_hooks/CMakeLists.txt index 94dfc8ab9f0..5b8f8bffb22 100644 --- a/core/device_hooks/CMakeLists.txt +++ b/core/device_hooks/CMakeLists.txt @@ -3,11 +3,28 @@ if(NOT GINKGO_BUILD_CUDA) $ cuda_hooks.cpp) ginkgo_compile_features(ginkgo_cuda) - target_link_libraries(ginkgo_cuda PUBLIC ginkgo_hip) ginkgo_default_includes(ginkgo_cuda) ginkgo_install_library(ginkgo_cuda cuda) endif() +if (NOT GINKGO_BUILD_DPCPP) + add_library(ginkgo_dpcpp + $ + dpcpp_hooks.cpp) + ginkgo_compile_features(ginkgo_dpcpp) + ginkgo_default_includes(ginkgo_dpcpp) + ginkgo_install_library(ginkgo_dpcpp dpcpp) +endif() + +if(NOT GINKGO_BUILD_HIP) + add_library(ginkgo_hip + $ + hip_hooks.cpp) + ginkgo_compile_features(ginkgo_hip) + ginkgo_default_includes(ginkgo_hip) + ginkgo_install_library(ginkgo_hip hip) +endif() + if (NOT GINKGO_BUILD_OMP) add_library(ginkgo_omp $ @@ -15,6 +32,7 @@ if (NOT GINKGO_BUILD_OMP) ginkgo_compile_features(ginkgo_omp) target_link_libraries(ginkgo_omp PUBLIC ginkgo_cuda) target_link_libraries(ginkgo_omp PUBLIC ginkgo_hip) + target_link_libraries(ginkgo_omp PUBLIC ginkgo_dpcpp) ginkgo_default_includes(ginkgo_omp) ginkgo_install_library(ginkgo_omp omp) endif() @@ -27,12 +45,3 @@ if (NOT GINKGO_BUILD_REFERENCE) ginkgo_default_includes(ginkgo_reference) ginkgo_install_library(ginkgo_reference reference) endif() - -if(NOT GINKGO_BUILD_HIP) - add_library(ginkgo_hip - $ - hip_hooks.cpp) - ginkgo_compile_features(ginkgo_hip) - ginkgo_default_includes(ginkgo_hip) - ginkgo_install_library(ginkgo_hip hip) -endif() diff --git a/core/device_hooks/cuda_hooks.cpp b/core/device_hooks/cuda_hooks.cpp index d41d77d24d9..5bd7f5cb09f 100644 --- a/core/device_hooks/cuda_hooks.cpp +++ b/core/device_hooks/cuda_hooks.cpp @@ -93,6 +93,11 @@ void CudaExecutor::raw_copy_to(const HipExecutor *, size_type num_bytes, GKO_NOT_COMPILED(cuda); +void CudaExecutor::raw_copy_to(const DpcppExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const + GKO_NOT_COMPILED(cuda); + + void CudaExecutor::synchronize() const GKO_NOT_COMPILED(cuda); diff --git a/core/device_hooks/dpcpp_hooks.cpp b/core/device_hooks/dpcpp_hooks.cpp new file mode 100644 index 00000000000..4c5a98800da --- /dev/null +++ b/core/device_hooks/dpcpp_hooks.cpp @@ -0,0 +1,118 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + + +#include +#include + + +#include +#include +#include +#include + + +namespace gko { + +version version_info::get_dpcpp_version() noexcept +{ + // We just return the version with a special "not compiled" tag in + // placeholder modules. + return {GKO_VERSION_STR, "not compiled"}; +} + + +std::shared_ptr DpcppExecutor::create( + int device_id, std::shared_ptr master, std::string device_type) +{ + return std::shared_ptr( + new DpcppExecutor(device_id, std::move(master), device_type)); +} + + +void OmpExecutor::raw_copy_to(const DpcppExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const + GKO_NOT_COMPILED(dpcpp); + + +void DpcppExecutor::raw_free(void *ptr) const noexcept +{ + // Free must never fail, as it can be called in destructors. + // If the nvidia module was not compiled, the library couldn't have + // allocated the memory, so there is no need to deallocate it. +} + + +void *DpcppExecutor::raw_alloc(size_type num_bytes) const + GKO_NOT_COMPILED(dpcpp); + + +void DpcppExecutor::raw_copy_to(const OmpExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const + GKO_NOT_COMPILED(dpcpp); + + +void DpcppExecutor::raw_copy_to(const CudaExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const + GKO_NOT_COMPILED(dpcpp); + +void DpcppExecutor::raw_copy_to(const HipExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const + GKO_NOT_COMPILED(dpcpp); + +void DpcppExecutor::raw_copy_to(const DpcppExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const + GKO_NOT_COMPILED(dpcpp); + + +void DpcppExecutor::synchronize() const GKO_NOT_COMPILED(dpcpp); + + +void DpcppExecutor::run(const Operation &op) const +{ + op.run(std::static_pointer_cast( + this->shared_from_this())); +} + + +int DpcppExecutor::get_num_devices(std::string) { return 0; } + + +void DpcppExecutor::set_gpu_property() {} + + +} // namespace gko + + +#define GKO_HOOK_MODULE dpcpp +#include "core/device_hooks/common_kernels.inc.cpp" +#undef GKO_HOOK_MODULE diff --git a/core/device_hooks/hip_hooks.cpp b/core/device_hooks/hip_hooks.cpp index a2e288b4157..a60b657bac7 100644 --- a/core/device_hooks/hip_hooks.cpp +++ b/core/device_hooks/hip_hooks.cpp @@ -90,6 +90,11 @@ void HipExecutor::raw_copy_to(const HipExecutor *, size_type num_bytes, GKO_NOT_COMPILED(hip); +void HipExecutor::raw_copy_to(const DpcppExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const + GKO_NOT_COMPILED(hip); + + void HipExecutor::synchronize() const GKO_NOT_COMPILED(hip); diff --git a/core/devices/CMakeLists.txt b/core/devices/CMakeLists.txt index 2a5626c0018..adfd93beb03 100644 --- a/core/devices/CMakeLists.txt +++ b/core/devices/CMakeLists.txt @@ -6,7 +6,8 @@ function(ginkgo_add_object_library name) set_target_properties(${name} PROPERTIES POSITION_INDEPENDENT_CODE ON) endfunction() -add_subdirectory(omp) add_subdirectory(cuda) +add_subdirectory(dpcpp) add_subdirectory(hip) +add_subdirectory(omp) add_subdirectory(reference) diff --git a/core/devices/dpcpp/CMakeLists.txt b/core/devices/dpcpp/CMakeLists.txt new file mode 100644 index 00000000000..f6fdb354ff2 --- /dev/null +++ b/core/devices/dpcpp/CMakeLists.txt @@ -0,0 +1,2 @@ +ginkgo_add_object_library(ginkgo_dpcpp_device + executor.cpp) diff --git a/core/devices/dpcpp/executor.cpp b/core/devices/dpcpp/executor.cpp new file mode 100644 index 00000000000..90b7eb67616 --- /dev/null +++ b/core/devices/dpcpp/executor.cpp @@ -0,0 +1,59 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include + + +#include +#include + + +#include +#include + + +namespace gko { + + +std::shared_ptr DpcppExecutor::get_master() noexcept +{ + return master_; +} + + +std::shared_ptr DpcppExecutor::get_master() const noexcept +{ + return master_; +} + + +} // namespace gko diff --git a/core/factorization/factorization_kernels.hpp b/core/factorization/factorization_kernels.hpp index f7c25964dde..96e87ed4362 100644 --- a/core/factorization/factorization_kernels.hpp +++ b/core/factorization/factorization_kernels.hpp @@ -132,6 +132,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace factorization { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace factorization +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/factorization/ilu_kernels.hpp b/core/factorization/ilu_kernels.hpp index 17602ac4ab4..1c71ceba897 100644 --- a/core/factorization/ilu_kernels.hpp +++ b/core/factorization/ilu_kernels.hpp @@ -95,6 +95,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace ilu_factorization { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace ilu_factorization +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/factorization/par_ict_kernels.hpp b/core/factorization/par_ict_kernels.hpp index f02b6ac7bb6..5cde5fa14e2 100644 --- a/core/factorization/par_ict_kernels.hpp +++ b/core/factorization/par_ict_kernels.hpp @@ -106,6 +106,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace par_ict_factorization { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace par_ict_factorization +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/factorization/par_ilu_kernels.hpp b/core/factorization/par_ilu_kernels.hpp index 09bc1dd2596..6d8451b32e7 100644 --- a/core/factorization/par_ilu_kernels.hpp +++ b/core/factorization/par_ilu_kernels.hpp @@ -97,6 +97,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace par_ilu_factorization { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace par_ilu_factorization +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/factorization/par_ilut_kernels.hpp b/core/factorization/par_ilut_kernels.hpp index 9bb19596c3f..714e797e553 100644 --- a/core/factorization/par_ilut_kernels.hpp +++ b/core/factorization/par_ilut_kernels.hpp @@ -143,6 +143,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace par_ilut_factorization { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace par_ilut_factorization +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/coo_kernels.hpp b/core/matrix/coo_kernels.hpp index 752798c4e2e..8ae0466db26 100644 --- a/core/matrix/coo_kernels.hpp +++ b/core/matrix/coo_kernels.hpp @@ -140,6 +140,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace coo { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace coo +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/csr_kernels.hpp b/core/matrix/csr_kernels.hpp index 92d9f462dfd..12abad81fa3 100644 --- a/core/matrix/csr_kernels.hpp +++ b/core/matrix/csr_kernels.hpp @@ -263,6 +263,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace csr { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace csr +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp index 30ad6ca4729..1609f142338 100644 --- a/core/matrix/dense_kernels.hpp +++ b/core/matrix/dense_kernels.hpp @@ -274,6 +274,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace dense { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace dense +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/diagonal_kernels.hpp b/core/matrix/diagonal_kernels.hpp index bb8d4c14e7d..556ab66d8a3 100644 --- a/core/matrix/diagonal_kernels.hpp +++ b/core/matrix/diagonal_kernels.hpp @@ -135,6 +135,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace diagonal { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace diagonal +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/ell_kernels.hpp b/core/matrix/ell_kernels.hpp index 049ed280b66..aa78e17881e 100644 --- a/core/matrix/ell_kernels.hpp +++ b/core/matrix/ell_kernels.hpp @@ -139,6 +139,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace ell { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace ell +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/hybrid_kernels.hpp b/core/matrix/hybrid_kernels.hpp index 788fe66e15b..fb813a0e428 100644 --- a/core/matrix/hybrid_kernels.hpp +++ b/core/matrix/hybrid_kernels.hpp @@ -104,6 +104,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace hybrid { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace hybrid +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/sellp_kernels.hpp b/core/matrix/sellp_kernels.hpp index 7b9b6fdf551..f57cf7087a8 100644 --- a/core/matrix/sellp_kernels.hpp +++ b/core/matrix/sellp_kernels.hpp @@ -130,6 +130,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace sellp { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace sellp +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/matrix/sparsity_csr_kernels.hpp b/core/matrix/sparsity_csr_kernels.hpp index 58ec58e789f..8d5e3570699 100644 --- a/core/matrix/sparsity_csr_kernels.hpp +++ b/core/matrix/sparsity_csr_kernels.hpp @@ -144,6 +144,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace sparsity_csr { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace sparsity_csr +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/preconditioner/isai_kernels.hpp b/core/preconditioner/isai_kernels.hpp index ce53d51cd3c..8093f454255 100644 --- a/core/preconditioner/isai_kernels.hpp +++ b/core/preconditioner/isai_kernels.hpp @@ -111,6 +111,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace isai { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace isai +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/preconditioner/jacobi_kernels.hpp b/core/preconditioner/jacobi_kernels.hpp index 12d232c26f8..96ae123df12 100644 --- a/core/preconditioner/jacobi_kernels.hpp +++ b/core/preconditioner/jacobi_kernels.hpp @@ -176,6 +176,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace jacobi { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace jacobi +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/bicg_kernels.hpp b/core/solver/bicg_kernels.hpp index 9ef21b3a243..c11c7f4e4a1 100644 --- a/core/solver/bicg_kernels.hpp +++ b/core/solver/bicg_kernels.hpp @@ -124,6 +124,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace bicg { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace bicg +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/bicgstab_kernels.hpp b/core/solver/bicgstab_kernels.hpp index 8b48151a50f..0ebd591a1fe 100644 --- a/core/solver/bicgstab_kernels.hpp +++ b/core/solver/bicgstab_kernels.hpp @@ -148,6 +148,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace bicgstab { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace bicgstab +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/cg_kernels.hpp b/core/solver/cg_kernels.hpp index 3a52974033a..1042b24cccd 100644 --- a/core/solver/cg_kernels.hpp +++ b/core/solver/cg_kernels.hpp @@ -119,6 +119,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace cg { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace cg +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/cgs_kernels.hpp b/core/solver/cgs_kernels.hpp index 1404303b2ce..e2622419573 100644 --- a/core/solver/cgs_kernels.hpp +++ b/core/solver/cgs_kernels.hpp @@ -135,6 +135,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace cgs { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace cgs +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/fcg_kernels.hpp b/core/solver/fcg_kernels.hpp index dc269f2fa19..89c3e1c3504 100644 --- a/core/solver/fcg_kernels.hpp +++ b/core/solver/fcg_kernels.hpp @@ -120,6 +120,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace fcg { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace fcg +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/gmres_kernels.hpp b/core/solver/gmres_kernels.hpp index 644a8cf708e..2aaa0fb509e 100644 --- a/core/solver/gmres_kernels.hpp +++ b/core/solver/gmres_kernels.hpp @@ -135,6 +135,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace gmres { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace gmres +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/ir_kernels.hpp b/core/solver/ir_kernels.hpp index 9fe59ba4a6c..ac76627755a 100644 --- a/core/solver/ir_kernels.hpp +++ b/core/solver/ir_kernels.hpp @@ -92,6 +92,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace ir { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace ir +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/lower_trs_kernels.hpp b/core/solver/lower_trs_kernels.hpp index 799c50129e0..4430de5b468 100644 --- a/core/solver/lower_trs_kernels.hpp +++ b/core/solver/lower_trs_kernels.hpp @@ -123,6 +123,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace lower_trs { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace lower_trs +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/solver/upper_trs_kernels.hpp b/core/solver/upper_trs_kernels.hpp index cce48ea2812..e49fe8cc8e7 100644 --- a/core/solver/upper_trs_kernels.hpp +++ b/core/solver/upper_trs_kernels.hpp @@ -123,6 +123,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace upper_trs { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace upper_trs +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES diff --git a/core/stop/criterion_kernels.hpp b/core/stop/criterion_kernels.hpp index 07eb8f2798c..c894e6337ca 100644 --- a/core/stop/criterion_kernels.hpp +++ b/core/stop/criterion_kernels.hpp @@ -87,6 +87,15 @@ GKO_DECLARE_SET_ALL_STATUSES_KERNEL(); } // namespace set_all_statuses } // namespace hip + + +namespace dpcpp { +namespace set_all_statuses { + +GKO_DECLARE_SET_ALL_STATUSES_KERNEL(); + +} // namespace set_all_statuses +} // namespace dpcpp } // namespace kernels } // namespace gko diff --git a/core/stop/residual_norm_kernels.hpp b/core/stop/residual_norm_kernels.hpp index 30407cf9b9f..94c9b7001ad 100644 --- a/core/stop/residual_norm_kernels.hpp +++ b/core/stop/residual_norm_kernels.hpp @@ -99,6 +99,15 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace hip +namespace dpcpp { +namespace residual_norm { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace residual_norm +} // namespace dpcpp + + #undef GKO_DECLARE_ALL_AS_TEMPLATES } // namespace kernels diff --git a/core/test/base/executor.cpp b/core/test/base/executor.cpp index 1b2e1b0698e..3a37115fedd 100644 --- a/core/test/base/executor.cpp +++ b/core/test/base/executor.cpp @@ -63,10 +63,14 @@ class ExampleOperation : public gko::Operation { { value = 3; } - void run(std::shared_ptr) const override + void run(std::shared_ptr) const override { value = 4; } + void run(std::shared_ptr) const override + { + value = 5; + } int &value; }; @@ -88,9 +92,10 @@ TEST(OmpExecutor, RunsCorrectLambdaOperation) auto omp_lambda = [&value]() { value = 1; }; auto cuda_lambda = [&value]() { value = 2; }; auto hip_lambda = [&value]() { value = 3; }; + auto dpcpp_lambda = [&value]() { value = 4; }; exec_ptr omp = gko::OmpExecutor::create(); - omp->run(omp_lambda, cuda_lambda, hip_lambda); + omp->run(omp_lambda, cuda_lambda, hip_lambda, dpcpp_lambda); ASSERT_EQ(1, value); } @@ -155,7 +160,7 @@ TEST(ReferenceExecutor, RunsCorrectOperation) exec_ptr ref = gko::ReferenceExecutor::create(); ref->run(ExampleOperation(value)); - ASSERT_EQ(4, value); + ASSERT_EQ(5, value); } @@ -165,9 +170,10 @@ TEST(ReferenceExecutor, RunsCorrectLambdaOperation) auto omp_lambda = [&value]() { value = 1; }; auto cuda_lambda = [&value]() { value = 2; }; auto hip_lambda = [&value]() { value = 3; }; + auto dpcpp_lambda = [&value]() { value = 4; }; exec_ptr ref = gko::ReferenceExecutor::create(); - ref->run(omp_lambda, cuda_lambda, hip_lambda); + ref->run(omp_lambda, cuda_lambda, hip_lambda, dpcpp_lambda); ASSERT_EQ(1, value); } @@ -289,10 +295,11 @@ TEST(CudaExecutor, RunsCorrectLambdaOperation) auto omp_lambda = [&value]() { value = 1; }; auto cuda_lambda = [&value]() { value = 2; }; auto hip_lambda = [&value]() { value = 3; }; + auto dpcpp_lambda = [&value]() { value = 4; }; exec_ptr cuda = gko::CudaExecutor::create(0, gko::OmpExecutor::create(), true); - cuda->run(omp_lambda, cuda_lambda, hip_lambda); + cuda->run(omp_lambda, cuda_lambda, hip_lambda, dpcpp_lambda); ASSERT_EQ(2, value); } @@ -360,9 +367,10 @@ TEST(HipExecutor, RunsCorrectLambdaOperation) auto omp_lambda = [&value]() { value = 1; }; auto cuda_lambda = [&value]() { value = 2; }; auto hip_lambda = [&value]() { value = 3; }; + auto dpcpp_lambda = [&value]() { value = 4; }; exec_ptr hip = gko::HipExecutor::create(0, gko::OmpExecutor::create()); - hip->run(omp_lambda, cuda_lambda, hip_lambda); + hip->run(omp_lambda, cuda_lambda, hip_lambda, dpcpp_lambda); ASSERT_EQ(3, value); } @@ -414,6 +422,50 @@ TEST(HipExecutor, CanSetDeviceResetBoolean) } +TEST(DpcppExecutor, RunsCorrectOperation) +{ + int value = 0; + exec_ptr dpcpp = + gko::DpcppExecutor::create(0, gko::OmpExecutor::create(), "all"); + + dpcpp->run(ExampleOperation(value)); + ASSERT_EQ(4, value); +} + + +TEST(DpcppExecutor, RunsCorrectLambdaOperation) +{ + int value = 0; + auto omp_lambda = [&value]() { value = 1; }; + auto cuda_lambda = [&value]() { value = 2; }; + auto hip_lambda = [&value]() { value = 3; }; + auto dpcpp_lambda = [&value]() { value = 4; }; + exec_ptr dpcpp = + gko::DpcppExecutor::create(0, gko::OmpExecutor::create(), "all"); + + dpcpp->run(omp_lambda, cuda_lambda, hip_lambda, dpcpp_lambda); + ASSERT_EQ(4, value); +} + + +TEST(DpcppExecutor, KnowsItsMaster) +{ + auto omp = gko::OmpExecutor::create(); + exec_ptr dpcpp = gko::DpcppExecutor::create(0, omp, "all"); + + ASSERT_EQ(omp, dpcpp->get_master()); +} + + +TEST(DpcppExecutor, KnowsItsDeviceId) +{ + auto omp = gko::OmpExecutor::create(); + auto dpcpp = gko::DpcppExecutor::create(0, omp, "all"); + + ASSERT_EQ(0, dpcpp->get_device_id()); +} + + template struct mock_free : T { /** diff --git a/cuda/CMakeLists.txt b/cuda/CMakeLists.txt index d3a73ab536c..ebba732320a 100644 --- a/cuda/CMakeLists.txt +++ b/cuda/CMakeLists.txt @@ -150,9 +150,6 @@ target_include_directories(ginkgo_cuda SYSTEM PRIVATE ${CUDA_INCLUDE_DIRS}) target_link_libraries(ginkgo_cuda PRIVATE ${CUDA_RUNTIME_LIBS} ${CUBLAS} ${CUSPARSE}) -# Need to link against ginkgo_hip for the `raw_copy_to(HipExecutor ...)` method -target_link_libraries(ginkgo_cuda PUBLIC ginkgo_hip) - cas_target_cuda_architectures(ginkgo_cuda ARCHITECTURES ${GINKGO_CUDA_ARCHITECTURES} UNSUPPORTED "20" "21") diff --git a/cuda/base/executor.cpp b/cuda/base/executor.cpp index 543e78131e0..7926ae98a21 100644 --- a/cuda/base/executor.cpp +++ b/cuda/base/executor.cpp @@ -127,31 +127,38 @@ void CudaExecutor::raw_copy_to(const OmpExecutor *, size_type num_bytes, } -void CudaExecutor::raw_copy_to(const CudaExecutor *dest, size_type num_bytes, +void CudaExecutor::raw_copy_to(const HipExecutor *dest, size_type num_bytes, const void *src_ptr, void *dest_ptr) const { +#if GINKGO_HIP_PLATFORM_NVCC == 1 if (num_bytes > 0) { cuda::device_guard g(this->get_device_id()); GKO_ASSERT_NO_CUDA_ERRORS( cudaMemcpyPeer(dest_ptr, dest->get_device_id(), src_ptr, this->get_device_id(), num_bytes)); } +#else + GKO_NOT_SUPPORTED(dest); +#endif } -void CudaExecutor::raw_copy_to(const HipExecutor *dest, size_type num_bytes, +void CudaExecutor::raw_copy_to(const DpcppExecutor *dest, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const +{ + GKO_NOT_SUPPORTED(dest); +} + + +void CudaExecutor::raw_copy_to(const CudaExecutor *dest, size_type num_bytes, const void *src_ptr, void *dest_ptr) const { -#if GINKGO_HIP_PLATFORM_NVCC == 1 if (num_bytes > 0) { cuda::device_guard g(this->get_device_id()); GKO_ASSERT_NO_CUDA_ERRORS( cudaMemcpyPeer(dest_ptr, dest->get_device_id(), src_ptr, this->get_device_id(), num_bytes)); } -#else - GKO_NOT_SUPPORTED(this); -#endif } diff --git a/cuda/test/base/cuda_executor.cu b/cuda/test/base/cuda_executor.cu index 2bcf5961bbd..7b944fd13a1 100644 --- a/cuda/test/base/cuda_executor.cu +++ b/cuda/test/base/cuda_executor.cu @@ -69,6 +69,11 @@ public: value = -3; } + void run(std::shared_ptr) const override + { + value = -4; + } + void run(std::shared_ptr) const override { cudaGetDevice(&value); @@ -239,7 +244,7 @@ TEST_F(CudaExecutor, CopiesDataFromCudaToCuda) omp->copy_from(cuda2.get(), 2, copy_cuda2, copy); EXPECT_EQ(3, copy[0]); ASSERT_EQ(8, copy[1]); - cuda->free(copy_cuda2); + cuda2->free(copy_cuda2); cuda->free(orig); } diff --git a/doc/headers/dpcpp_executor.hpp b/doc/headers/dpcpp_executor.hpp new file mode 100644 index 00000000000..83df1f5b292 --- /dev/null +++ b/doc/headers/dpcpp_executor.hpp @@ -0,0 +1,40 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +/** + * @defgroup exec_omp OpenMP Executor + * + * @brief A module dedicated to the implementation and usage of the OpenMP + * executor in Ginkgo. + * + * @ingroup Executor + */ diff --git a/dpcpp/CMakeLists.txt b/dpcpp/CMakeLists.txt new file mode 100644 index 00000000000..55401c838fe --- /dev/null +++ b/dpcpp/CMakeLists.txt @@ -0,0 +1,57 @@ +if (NOT GKO_CAN_COMPILE_DPCPP) + message(FATAL_ERROR "The CMAKE_CXX_COMPILER compiler, which is " + "${CMAKE_CXX_COMPILER} cannot compile DPC++ code!") +endif() + +add_library(ginkgo_dpcpp $ "") +target_sources(ginkgo_dpcpp + PRIVATE + base/version.dp.cpp + base/executor.dp.cpp + components/absolute_array.dp.cpp + components/fill_array.dp.cpp + components/precision_conversion.dp.cpp + components/prefix_sum.dp.cpp + factorization/ilu_kernels.dp.cpp + factorization/factorization_kernels.dp.cpp + factorization/par_ict_kernels.dp.cpp + factorization/par_ilu_kernels.dp.cpp + factorization/par_ilut_kernels.dp.cpp + matrix/coo_kernels.dp.cpp + matrix/csr_kernels.dp.cpp + matrix/dense_kernels.dp.cpp + matrix/diagonal_kernels.dp.cpp + matrix/ell_kernels.dp.cpp + matrix/hybrid_kernels.dp.cpp + matrix/sellp_kernels.dp.cpp + matrix/sparsity_csr_kernels.dp.cpp + preconditioner/isai_kernels.dp.cpp + preconditioner/jacobi_kernels.dp.cpp + solver/bicg_kernels.dp.cpp + solver/bicgstab_kernels.dp.cpp + solver/cg_kernels.dp.cpp + solver/cgs_kernels.dp.cpp + solver/fcg_kernels.dp.cpp + solver/gmres_kernels.dp.cpp + solver/ir_kernels.dp.cpp + solver/lower_trs_kernels.dp.cpp + solver/upper_trs_kernels.dp.cpp + stop/criterion_kernels.dp.cpp + stop/residual_norm_kernels.dp.cpp) + +ginkgo_compile_features(ginkgo_dpcpp) + +set(GINKGO_DPCPP_FLAGS ${GINKGO_COMPILER_FLAGS} -fsycl) +target_compile_options(ginkgo_dpcpp PRIVATE "${GINKGO_DPCPP_FLAGS}") +target_compile_features(ginkgo_dpcpp PRIVATE cxx_std_17) + +ginkgo_default_includes(ginkgo_dpcpp) +ginkgo_install_library(ginkgo_dpcpp dpcpp) + +if (GINKGO_CHECK_CIRCULAR_DEPS) + ginkgo_check_headers(ginkgo_dpcpp) +endif() + +if(GINKGO_BUILD_TESTS) + add_subdirectory(test) +endif() diff --git a/dpcpp/base/executor.dp.cpp b/dpcpp/base/executor.dp.cpp new file mode 100644 index 00000000000..05db4fda198 --- /dev/null +++ b/dpcpp/base/executor.dp.cpp @@ -0,0 +1,181 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include + + +#include +#include +#include + + +#include + + +#include +#include + + +namespace gko { +namespace detail { + + +const std::vector get_devices(std::string device_type) +{ + std::map device_type_map{ + {"accelerator", sycl::info::device_type::accelerator}, + {"all", sycl::info::device_type::all}, + {"cpu", sycl::info::device_type::cpu}, + {"gpu", sycl::info::device_type::gpu}}; + return sycl::device::get_devices(device_type_map.at(device_type)); +} + + +} // namespace detail + + +void OmpExecutor::raw_copy_to(const DpcppExecutor *dest, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const +{ + if (num_bytes > 0) { + dest->get_queue()->memcpy(dest_ptr, src_ptr, num_bytes).wait(); + } +} + + +std::shared_ptr DpcppExecutor::create( + int device_id, std::shared_ptr master, std::string device_type) +{ + return std::shared_ptr( + new DpcppExecutor(device_id, std::move(master), device_type)); +} + + +void DpcppExecutor::raw_free(void *ptr) const noexcept +{ + sycl::free(ptr, queue_->get_context()); +} + + +void *DpcppExecutor::raw_alloc(size_type num_bytes) const +{ + void *dev_ptr = sycl::malloc_device(num_bytes, *queue_.get()); + GKO_ENSURE_ALLOCATED(dev_ptr, "DPC++", num_bytes); + return dev_ptr; +} + + +void DpcppExecutor::raw_copy_to(const OmpExecutor *, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const +{ + if (num_bytes > 0) { + queue_->memcpy(dest_ptr, src_ptr, num_bytes).wait(); + } +} + + +void DpcppExecutor::raw_copy_to(const CudaExecutor *dest, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const +{ + // TODO: later when possible, if we have DPC++ with a CUDA backend + // support/compiler, we could maybe support native copies? + GKO_NOT_SUPPORTED(dest); +} + + +void DpcppExecutor::raw_copy_to(const HipExecutor *dest, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const +{ + GKO_NOT_SUPPORTED(dest); +} + + +void DpcppExecutor::raw_copy_to(const DpcppExecutor *dest, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const +{ + if (num_bytes > 0) { + // TODO: does this work? Or is it needed to go through host? + dest->get_queue()->memcpy(dest_ptr, src_ptr, num_bytes).wait(); + } +} + + +void DpcppExecutor::synchronize() const { queue_->wait_and_throw(); } + + +void DpcppExecutor::run(const Operation &op) const +{ + this->template log(this, &op); + op.run(std::static_pointer_cast( + this->shared_from_this())); + this->template log(this, &op); +} + + +int DpcppExecutor::get_num_devices(std::string device_type) +{ + return detail::get_devices(device_type).size(); +} + + +void delete_queue(sycl::queue *queue) +{ + queue->wait(); + delete queue; +} + + +void DpcppExecutor::set_gpu_property() +{ + auto device = detail::get_devices(device_type_)[device_id_]; + try { + subgroup_sizes_ = + device.get_info(); + } catch (cl::sycl::runtime_error &err) { + GKO_NOT_SUPPORTED(device); + } + num_computing_units_ = + device.get_info(); + auto max_workitem_sizes = + device.get_info(); + // There is no way to get the dimension of a sycl::id object + for (std::size_t i = 0; i < 3; i++) { + max_workitem_sizes_.push_back(max_workitem_sizes[i]); + } + max_workgroup_size_ = + device.get_info(); + auto *queue = new sycl::queue{device, sycl::property::queue::in_order{}}; + queue_ = std::move(queue_manager{queue, delete_queue}); +} + + +} // namespace gko diff --git a/dpcpp/base/version.dp.cpp b/dpcpp/base/version.dp.cpp new file mode 100644 index 00000000000..229a72ca72c --- /dev/null +++ b/dpcpp/base/version.dp.cpp @@ -0,0 +1,48 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include + + +namespace gko { + + +version version_info::get_dpcpp_version() noexcept +{ + // When compiling the module, the header version is the same as the library + // version. Mismatch between the header and the module versions may happen + // if using shared libraries from different versions of Ginkgo. + return version_info::get_header_version(); +} + + +} // namespace gko diff --git a/dpcpp/components/absolute_array.dp.cpp b/dpcpp/components/absolute_array.dp.cpp new file mode 100644 index 00000000000..9603dacbc8b --- /dev/null +++ b/dpcpp/components/absolute_array.dp.cpp @@ -0,0 +1,82 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/components/absolute_array.hpp" + + +#include + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +namespace components { + + +template +void inplace_absolute_array(std::shared_ptr exec, + ValueType *data, size_type n) +{ + exec->get_queue()->submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx_id) { + const int idx = idx_id[0]; + data[idx] = dpcpp::abs(data[idx]); + }); + }); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL); + + +template +void outplace_absolute_array(std::shared_ptr exec, + const ValueType *in, size_type n, + remove_complex *out) +{ + exec->get_queue()->submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx_id) { + const int idx = idx_id[0]; + out[idx] = dpcpp::abs(in[idx]); + }); + }); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL); + + +} // namespace components +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/components/csr_spgeam.dp.hpp b/dpcpp/components/csr_spgeam.dp.hpp new file mode 100644 index 00000000000..77303d7dd7a --- /dev/null +++ b/dpcpp/components/csr_spgeam.dp.hpp @@ -0,0 +1,76 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#ifndef GKO_DPCPP_COMPONENTS_CSR_SPGEAM_HPP_ +#define GKO_DPCPP_COMPONENTS_CSR_SPGEAM_HPP_ + + +#include + + +#include + + +#include "core/base/utils.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { + + +/** + * Adds two (sorted) sparse matrices. + * + * Calls begin_cb(row) on each row to initialize row-local data + * Calls entry_cb(row, col, a_val, b_val, local_data) on each output non-zero + * Calls end_cb(row, local_data) on each row to finalize row-local data + * + * If the three functions are thread-safe, the whole invocation is. + */ +template +void abstract_spgeam(const matrix::Csr *a, + const matrix::Csr *b, + BeginCallback begin_cb, EntryCallback entry_cb, + EndCallback end_cb) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace dpcpp +} // namespace kernels +} // namespace gko + + +#endif // GKO_DPCPP_COMPONENTS_CSR_SPGEAM_HPP_ diff --git a/dpcpp/components/fill_array.dp.cpp b/dpcpp/components/fill_array.dp.cpp new file mode 100644 index 00000000000..f8c260cdab6 --- /dev/null +++ b/dpcpp/components/fill_array.dp.cpp @@ -0,0 +1,65 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/components/fill_array.hpp" + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +namespace components { + + +template +void fill_array(std::shared_ptr exec, ValueType *array, + size_type n, ValueType val) +{ + exec->get_queue()->submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx_id) { + const int idx = idx_id[0]; + array[idx] = val; + }); + }); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_FILL_ARRAY_KERNEL); +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE(GKO_DECLARE_FILL_ARRAY_KERNEL); +template GKO_DECLARE_FILL_ARRAY_KERNEL(size_type); + + +} // namespace components +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/components/format_conversion.dp.hpp b/dpcpp/components/format_conversion.dp.hpp new file mode 100644 index 00000000000..f1dba5eed78 --- /dev/null +++ b/dpcpp/components/format_conversion.dp.hpp @@ -0,0 +1,91 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include +#include + + +#include + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { + + +/** + * @internal + * + * Converts an array of indexes `idxs` in any order to an array of pointers + * `ptrs`. This is used for transposing a csr matrix when calculating the row + * pointers of the transposed matrix out of the column indices of the original + * matrix. + */ +template +inline void convert_unsorted_idxs_to_ptrs(const IndexType *idxs, + size_type num_nonzeros, + IndexType *ptrs, size_type length) +{ + GKO_NOT_IMPLEMENTED; +} + + +/** + * @internal + * + * Converts an array of indexes `idxs` which are already stored in an increasing + * order to an array of pointers `ptrs`. This is used to calculate the row + * pointers when converting a coo matrix to a csr matrix. + */ +template +inline void convert_sorted_idxs_to_ptrs(const IndexType *idxs, + size_type num_nonzeros, IndexType *ptrs, + size_type length) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +inline void convert_ptrs_to_idxs(const IndexType *ptrs, size_type num_rows, + IndexType *idxs) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/components/matrix_operations.dp.hpp b/dpcpp/components/matrix_operations.dp.hpp new file mode 100644 index 00000000000..4efb2d839f2 --- /dev/null +++ b/dpcpp/components/matrix_operations.dp.hpp @@ -0,0 +1,65 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#ifndef GKO_DPCPP_COMPONENTS_MATRIX_OPERATIONS_HPP_ +#define GKO_DPCPP_COMPONENTS_MATRIX_OPERATIONS_HPP_ + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { + + +/** + * @internal + * + * Computes the infinity norm of a column-major matrix. + */ +template +remove_complex compute_inf_norm(size_type num_rows, + size_type num_cols, + const ValueType *matrix, + size_type stride) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace dpcpp +} // namespace kernels +} // namespace gko + + +#endif // GKO_DPCPP_COMPONENTS_MATRIX_OPERATIONS_HPP_ diff --git a/dpcpp/components/precision_conversion.dp.cpp b/dpcpp/components/precision_conversion.dp.cpp new file mode 100644 index 00000000000..c508e853ad6 --- /dev/null +++ b/dpcpp/components/precision_conversion.dp.cpp @@ -0,0 +1,63 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/components/precision_conversion.hpp" + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +namespace components { + + +template +void convert_precision(std::shared_ptr exec, + size_type size, const SourceType *in, TargetType *out) +{ + exec->get_queue()->submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<1>{size}, [=](sycl::id<1> idx_id) { + const int idx = idx_id[0]; + out[idx] = in[idx]; + }); + }); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_CONVERSION(GKO_DECLARE_CONVERT_PRECISION_KERNEL); + + +} // namespace components +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/components/prefix_sum.dp.cpp b/dpcpp/components/prefix_sum.dp.cpp new file mode 100644 index 00000000000..21ca279b696 --- /dev/null +++ b/dpcpp/components/prefix_sum.dp.cpp @@ -0,0 +1,64 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/components/prefix_sum.hpp" + + +#include + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +namespace components { + + +template +void prefix_sum(std::shared_ptr exec, IndexType *counts, + size_type num_entries) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE(GKO_DECLARE_PREFIX_SUM_KERNEL); + +// instantiate for size_type as well, as this is used in the Sellp format +template GKO_DECLARE_PREFIX_SUM_KERNEL(size_type); + + +} // namespace components +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/factorization/factorization_kernels.dp.cpp b/dpcpp/factorization/factorization_kernels.dp.cpp new file mode 100644 index 00000000000..5a571b86ecd --- /dev/null +++ b/dpcpp/factorization/factorization_kernels.dp.cpp @@ -0,0 +1,174 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/factorization/factorization_kernels.hpp" + + +#include +#include + + +#include +#include + + +#include "core/components/prefix_sum.hpp" +#include "core/matrix/csr_builder.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The factorization namespace. + * + * @ingroup factor + */ +namespace factorization { + + +namespace kernel { +namespace detail { + + +template +struct find_helper { + template + static inline bool find(ForwardIt first, ForwardIt last, IndexType value) + { + return std::find(first, last, value) != last; + } +}; + + +template <> +struct find_helper { + template + static inline bool find(ForwardIt first, ForwardIt last, IndexType value) + { + return std::binary_search(first, last, value); + } +}; + + +} // namespace detail + + +template +void find_missing_diagonal_elements( + const matrix::Csr *mtx, + IndexType *elements_to_add_per_row, bool *changes_required) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void add_missing_diagonal_elements(const matrix::Csr *mtx, + ValueType *new_values, + IndexType *new_col_idxs, + const IndexType *row_ptrs_addition) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace kernel + + +template +void add_diagonal_elements(std::shared_ptr exec, + matrix::Csr *mtx, + bool is_sorted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_FACTORIZATION_ADD_DIAGONAL_ELEMENTS_KERNEL); + + +template +void initialize_row_ptrs_l_u( + std::shared_ptr exec, + const matrix::Csr *system_matrix, + IndexType *l_row_ptrs, IndexType *u_row_ptrs) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_FACTORIZATION_INITIALIZE_ROW_PTRS_L_U_KERNEL); + + +template +void initialize_l_u(std::shared_ptr exec, + const matrix::Csr *system_matrix, + matrix::Csr *csr_l, + matrix::Csr *csr_u) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_FACTORIZATION_INITIALIZE_L_U_KERNEL); + + +template +void initialize_row_ptrs_l( + std::shared_ptr exec, + const matrix::Csr *system_matrix, + IndexType *l_row_ptrs) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_FACTORIZATION_INITIALIZE_ROW_PTRS_L_KERNEL); + + +template +void initialize_l(std::shared_ptr exec, + const matrix::Csr *system_matrix, + matrix::Csr *csr_l, bool diag_sqrt) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_FACTORIZATION_INITIALIZE_L_KERNEL); + + +} // namespace factorization +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/factorization/ilu_kernels.dp.cpp b/dpcpp/factorization/ilu_kernels.dp.cpp new file mode 100644 index 00000000000..1d854a35a3d --- /dev/null +++ b/dpcpp/factorization/ilu_kernels.dp.cpp @@ -0,0 +1,58 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/factorization/ilu_kernels.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The ilu factorization namespace. + * + * @ingroup factor + */ +namespace ilu_factorization { + + +template +void compute_lu(std::shared_ptr exec, + matrix::Csr *m) GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ILU_COMPUTE_LU_KERNEL); + + +} // namespace ilu_factorization +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/factorization/par_ict_kernels.dp.cpp b/dpcpp/factorization/par_ict_kernels.dp.cpp new file mode 100644 index 00000000000..e11b55244cd --- /dev/null +++ b/dpcpp/factorization/par_ict_kernels.dp.cpp @@ -0,0 +1,95 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/factorization/par_ict_kernels.hpp" + + +#include +#include +#include +#include + + +#include +#include +#include +#include + + +#include "core/base/utils.hpp" +#include "core/components/prefix_sum.hpp" +#include "core/matrix/csr_builder.hpp" +#include "dpcpp/components/csr_spgeam.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The parallel ICT factorization namespace. + * + * @ingroup factor + */ +namespace par_ict_factorization { + + +template +void compute_factor(std::shared_ptr exec, + const matrix::Csr *a, + matrix::Csr *l, + const matrix::Coo *) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ICT_COMPUTE_FACTOR_KERNEL); + + +template +void add_candidates(std::shared_ptr exec, + const matrix::Csr *llt, + const matrix::Csr *a, + const matrix::Csr *l, + matrix::Csr *l_new) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ICT_ADD_CANDIDATES_KERNEL); + + +} // namespace par_ict_factorization +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/factorization/par_ilu_kernels.dp.cpp b/dpcpp/factorization/par_ilu_kernels.dp.cpp new file mode 100644 index 00000000000..baa8b4a4481 --- /dev/null +++ b/dpcpp/factorization/par_ilu_kernels.dp.cpp @@ -0,0 +1,72 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/factorization/par_ilu_kernels.hpp" + + +#include + + +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The parallel ILU factorization namespace. + * + * @ingroup factor + */ +namespace par_ilu_factorization { + + +template +void compute_l_u_factors(std::shared_ptr exec, + size_type iterations, + const matrix::Coo *system_matrix, + matrix::Csr *l_factor, + matrix::Csr *u_factor) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ILU_COMPUTE_L_U_FACTORS_KERNEL); + + +} // namespace par_ilu_factorization +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/factorization/par_ilut_kernels.dp.cpp b/dpcpp/factorization/par_ilut_kernels.dp.cpp new file mode 100644 index 00000000000..fd26a1a7dfa --- /dev/null +++ b/dpcpp/factorization/par_ilut_kernels.dp.cpp @@ -0,0 +1,168 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/factorization/par_ilut_kernels.hpp" + + +#include +#include +#include +#include + + +#include + + +#include +#include +#include +#include + + +#include "core/base/utils.hpp" +#include "core/components/prefix_sum.hpp" +#include "core/matrix/coo_builder.hpp" +#include "core/matrix/csr_builder.hpp" +#include "dpcpp/components/csr_spgeam.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The parallel ILUT factorization namespace. + * + * @ingroup factor + */ +namespace par_ilut_factorization { + + +template +void threshold_select(std::shared_ptr exec, + const matrix::Csr *m, + IndexType rank, Array &tmp, + Array> &, + remove_complex &threshold) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ILUT_THRESHOLD_SELECT_KERNEL); + + +/** + * Removes all the elements from the input matrix for which pred is false. + * Stores the result in m_out and (if non-null) m_out_coo. + * pred(row, nz) is called for each entry, where nz is the index in + * values/col_idxs. + */ +template +void abstract_filter(std::shared_ptr exec, + const matrix::Csr *m, + matrix::Csr *m_out, + matrix::Coo *m_out_coo, + Predicate pred) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void threshold_filter(std::shared_ptr exec, + const matrix::Csr *m, + remove_complex threshold, + matrix::Csr *m_out, + matrix::Coo *m_out_coo, bool) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ILUT_THRESHOLD_FILTER_KERNEL); + + +constexpr auto bucket_count = 1 << sampleselect_searchtree_height; +constexpr auto sample_size = bucket_count * sampleselect_oversampling; + + +template +void threshold_filter_approx(std::shared_ptr exec, + const matrix::Csr *m, + IndexType rank, Array &tmp, + remove_complex &threshold, + matrix::Csr *m_out, + matrix::Coo *m_out_coo) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ILUT_THRESHOLD_FILTER_APPROX_KERNEL); + + +template +void compute_l_u_factors(std::shared_ptr exec, + const matrix::Csr *a, + matrix::Csr *l, + const matrix::Coo *, + matrix::Csr *u, + const matrix::Coo *, + matrix::Csr *u_csc) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ILUT_COMPUTE_LU_FACTORS_KERNEL); + + +template +void add_candidates(std::shared_ptr exec, + const matrix::Csr *lu, + const matrix::Csr *a, + const matrix::Csr *l, + const matrix::Csr *u, + matrix::Csr *l_new, + matrix::Csr *u_new) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_PAR_ILUT_ADD_CANDIDATES_KERNEL); + + +} // namespace par_ilut_factorization +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/get_info.cmake b/dpcpp/get_info.cmake new file mode 100644 index 00000000000..475a53f5dcf --- /dev/null +++ b/dpcpp/get_info.cmake @@ -0,0 +1,4 @@ +ginkgo_print_module_header(${detailed_log} "DPCPP") +ginkgo_print_module_footer(${detailed_log} "DPCPP variables:") +ginkgo_print_variable(${detailed_log} "GINKGO_DPCPP_FLAGS") +ginkgo_print_module_footer(${detailed_log} "") diff --git a/dpcpp/matrix/coo_kernels.dp.cpp b/dpcpp/matrix/coo_kernels.dp.cpp new file mode 100644 index 00000000000..acadcc41c23 --- /dev/null +++ b/dpcpp/matrix/coo_kernels.dp.cpp @@ -0,0 +1,163 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/coo_kernels.hpp" + + +#include + + +#include +#include +#include +#include + + +#include "dpcpp/components/format_conversion.hpp" + + +namespace gko { +namespace kernels { +/** + * @brief DPCPP namespace. + * + * @ingroup dpcpp + */ +namespace dpcpp { +/** + * @brief The Coordinate matrix format namespace. + * + * @ingroup coo + */ +namespace coo { + + +template +void spmv(std::shared_ptr exec, + const matrix::Coo *a, + const matrix::Dense *b, matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_SPMV_KERNEL); + + +template +void advanced_spmv(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Coo *a, + const matrix::Dense *b, + const matrix::Dense *beta, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_COO_ADVANCED_SPMV_KERNEL); + + +template +void spmv2(std::shared_ptr exec, + const matrix::Coo *a, + const matrix::Dense *b, matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_SPMV2_KERNEL); + + +template +void advanced_spmv2(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Coo *a, + const matrix::Dense *b, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_COO_ADVANCED_SPMV2_KERNEL); + + +template +void convert_row_idxs_to_ptrs(std::shared_ptr exec, + const IndexType *idxs, size_type num_nonzeros, + IndexType *ptrs, size_type length) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void convert_to_csr(std::shared_ptr exec, + const matrix::Coo *source, + matrix::Csr *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_COO_CONVERT_TO_CSR_KERNEL); + + +template +void convert_to_dense(std::shared_ptr exec, + const matrix::Coo *source, + matrix::Dense *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_COO_CONVERT_TO_DENSE_KERNEL); + + +template +void extract_diagonal(std::shared_ptr exec, + const matrix::Coo *orig, + matrix::Diagonal *diag) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_COO_EXTRACT_DIAGONAL_KERNEL); + + +} // namespace coo +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/csr_kernels.dp.cpp b/dpcpp/matrix/csr_kernels.dp.cpp new file mode 100644 index 00000000000..5a80cb12eaa --- /dev/null +++ b/dpcpp/matrix/csr_kernels.dp.cpp @@ -0,0 +1,425 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/csr_kernels.hpp" + + +#include +#include +#include + + +#include + + +#include +#include +#include +#include +#include +#include + + +#include "core/base/allocator.hpp" +#include "core/base/iterator_factory.hpp" +#include "core/components/prefix_sum.hpp" +#include "core/matrix/csr_builder.hpp" +#include "dpcpp/components/csr_spgeam.hpp" +#include "dpcpp/components/format_conversion.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Compressed sparse row matrix format namespace. + * + * @ingroup csr + */ +namespace csr { + + +template +void spmv(std::shared_ptr exec, + const matrix::Csr *a, + const matrix::Dense *b, matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_SPMV_KERNEL); + + +template +void advanced_spmv(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Csr *a, + const matrix::Dense *b, + const matrix::Dense *beta, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_ADVANCED_SPMV_KERNEL); + + +template +void spgemm_insert_row(unordered_set &cols, + const matrix::Csr *c, + size_type row) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void spgemm_insert_row2(unordered_set &cols, + const matrix::Csr *a, + const matrix::Csr *b, + size_type row) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void spgemm_accumulate_row(map &cols, + const matrix::Csr *c, + ValueType scale, size_type row) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void spgemm_accumulate_row2(map &cols, + const matrix::Csr *a, + const matrix::Csr *b, + ValueType scale, size_type row) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void spgemm(std::shared_ptr exec, + const matrix::Csr *a, + const matrix::Csr *b, + matrix::Csr *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_SPGEMM_KERNEL); + + +template +void advanced_spgemm(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Csr *a, + const matrix::Csr *b, + const matrix::Dense *beta, + const matrix::Csr *d, + matrix::Csr *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_ADVANCED_SPGEMM_KERNEL); + + +template +void spgeam(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Csr *a, + const matrix::Dense *beta, + const matrix::Csr *b, + matrix::Csr *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_SPGEAM_KERNEL); + + +template +void convert_row_ptrs_to_idxs(std::shared_ptr exec, + const IndexType *ptrs, size_type num_rows, + IndexType *idxs) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void convert_to_coo(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Coo *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CONVERT_TO_COO_KERNEL); + + +template +void convert_to_dense(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Dense *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CONVERT_TO_DENSE_KERNEL); + + +template +void convert_to_sellp(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Sellp *result) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CONVERT_TO_SELLP_KERNEL); + + +template +void convert_to_ell(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Ell *result) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CONVERT_TO_ELL_KERNEL); + + +template +inline void convert_csr_to_csc(size_type num_rows, const IndexType *row_ptrs, + const IndexType *col_idxs, + const ValueType *csr_vals, IndexType *row_idxs, + IndexType *col_ptrs, ValueType *csc_vals, + UnaryOperator op) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void transpose_and_transform(std::shared_ptr exec, + matrix::Csr *trans, + const matrix::Csr *orig, + UnaryOperator op) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void transpose(std::shared_ptr exec, + const matrix::Csr *orig, + matrix::Csr *trans) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_TRANSPOSE_KERNEL); + + +template +void conj_transpose(std::shared_ptr exec, + const matrix::Csr *orig, + matrix::Csr *trans) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CONJ_TRANSPOSE_KERNEL); + + +template +void calculate_total_cols(std::shared_ptr exec, + const matrix::Csr *source, + size_type *result, size_type stride_factor, + size_type slice_size) GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CALCULATE_TOTAL_COLS_KERNEL); + + +template +void calculate_max_nnz_per_row(std::shared_ptr exec, + const matrix::Csr *source, + size_type *result) GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CALCULATE_MAX_NNZ_PER_ROW_KERNEL); + + +template +void convert_to_hybrid(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Hybrid *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CONVERT_TO_HYBRID_KERNEL); + + +template +void row_permute_impl(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Csr *orig, + matrix::Csr *row_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void row_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Csr *orig, + matrix::Csr *row_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_ROW_PERMUTE_KERNEL); + + +template +void inverse_row_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Csr *orig, + matrix::Csr *row_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_INVERSE_ROW_PERMUTE_KERNEL); + + +template +void column_permute_impl(const Array *permutation_indices, + const matrix::Csr *orig, + matrix::Csr *column_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void column_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Csr *orig, + matrix::Csr *column_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_COLUMN_PERMUTE_KERNEL); + + +template +void inverse_column_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Csr *orig, + matrix::Csr *column_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_INVERSE_COLUMN_PERMUTE_KERNEL); + + +template +void calculate_nonzeros_per_row(std::shared_ptr exec, + const matrix::Csr *source, + Array *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_CALCULATE_NONZEROS_PER_ROW_KERNEL); + + +template +void sort_by_column_index(std::shared_ptr exec, + matrix::Csr *to_sort) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_SORT_BY_COLUMN_INDEX); + + +template +void is_sorted_by_column_index( + std::shared_ptr exec, + const matrix::Csr *to_check, bool *is_sorted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX); + + +template +void extract_diagonal(std::shared_ptr exec, + const matrix::Csr *orig, + matrix::Diagonal *diag) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_EXTRACT_DIAGONAL); + + +} // namespace csr +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/dense_kernels.dp.cpp b/dpcpp/matrix/dense_kernels.dp.cpp new file mode 100644 index 00000000000..5f76e1b5b97 --- /dev/null +++ b/dpcpp/matrix/dense_kernels.dp.cpp @@ -0,0 +1,375 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/dense_kernels.hpp" + + +#include + + +#include + + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + + +#include "core/components/prefix_sum.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Dense matrix format namespace. + * + * @ingroup dense + */ +namespace dense { + + +template +void simple_apply(std::shared_ptr exec, + const matrix::Dense *a, + const matrix::Dense *b, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_SIMPLE_APPLY_KERNEL); + + +template +void apply(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Dense *a, const matrix::Dense *b, + const matrix::Dense *beta, matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_APPLY_KERNEL); + + +template +void scale(std::shared_ptr exec, + const matrix::Dense *alpha, matrix::Dense *x) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_SCALE_KERNEL); + + +template +void add_scaled(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Dense *x, matrix::Dense *y) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_ADD_SCALED_KERNEL); + + +template +void add_scaled_diag(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Diagonal *x, + matrix::Dense *y) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_ADD_SCALED_DIAG_KERNEL); + + +template +void compute_dot(std::shared_ptr exec, + const matrix::Dense *x, + const matrix::Dense *y, + matrix::Dense *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_COMPUTE_DOT_KERNEL); + + +template +void compute_norm2(std::shared_ptr exec, + const matrix::Dense *x, + matrix::Dense> *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_COMPUTE_NORM2_KERNEL); + + +template +void convert_to_coo(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Coo *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_COO_KERNEL); + + +template +void convert_to_csr(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Csr *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL); + + +template +void convert_to_ell(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Ell *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL); + + +template +void convert_to_hybrid(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Hybrid *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL); + + +template +void convert_to_sellp(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Sellp *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); + + +template +void convert_to_sparsity_csr(std::shared_ptr exec, + const matrix::Dense *source, + matrix::SparsityCsr *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL); + + +template +void count_nonzeros(std::shared_ptr exec, + const matrix::Dense *source, size_type *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL); + + +template +void calculate_max_nnz_per_row(std::shared_ptr exec, + const matrix::Dense *source, + size_type *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE( + GKO_DECLARE_DENSE_CALCULATE_MAX_NNZ_PER_ROW_KERNEL); + + +template +void calculate_nonzeros_per_row(std::shared_ptr exec, + const matrix::Dense *source, + Array *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE( + GKO_DECLARE_DENSE_CALCULATE_NONZEROS_PER_ROW_KERNEL); + + +template +void calculate_total_cols(std::shared_ptr exec, + const matrix::Dense *source, + size_type *result, size_type stride_factor, + size_type slice_size) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE( + GKO_DECLARE_DENSE_CALCULATE_TOTAL_COLS_KERNEL); + + +template +void transpose(std::shared_ptr exec, + const matrix::Dense *orig, + matrix::Dense *trans) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_TRANSPOSE_KERNEL); + + +template +void conj_transpose(std::shared_ptr exec, + const matrix::Dense *orig, + matrix::Dense *trans) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CONJ_TRANSPOSE_KERNEL); + + +template +void row_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Dense *orig, + matrix::Dense *row_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_ROW_PERMUTE_KERNEL); + + +template +void column_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Dense *orig, + matrix::Dense *column_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_COLUMN_PERMUTE_KERNEL); + + +template +void inverse_row_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Dense *orig, + matrix::Dense *row_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_INVERSE_ROW_PERMUTE_KERNEL); + + +template +void inverse_column_permute(std::shared_ptr exec, + const Array *permutation_indices, + const matrix::Dense *orig, + matrix::Dense *column_permuted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_INVERSE_COLUMN_PERMUTE_KERNEL); + + +template +void extract_diagonal(std::shared_ptr exec, + const matrix::Dense *orig, + matrix::Diagonal *diag) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_EXTRACT_DIAGONAL_KERNEL); + + +template +void inplace_absolute_dense(std::shared_ptr exec, + matrix::Dense *source) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL); + + +template +void outplace_absolute_dense(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); + + +} // namespace dense +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/diagonal_kernels.dp.cpp b/dpcpp/matrix/diagonal_kernels.dp.cpp new file mode 100644 index 00000000000..c2090a17dc4 --- /dev/null +++ b/dpcpp/matrix/diagonal_kernels.dp.cpp @@ -0,0 +1,131 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/diagonal_kernels.hpp" + + +#include + + +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Diagonal matrix format namespace. + * + * @ingroup diagonal + */ +namespace diagonal { + + +template +void apply_to_dense(std::shared_ptr exec, + const matrix::Diagonal *a, + const matrix::Dense *b, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DIAGONAL_APPLY_TO_DENSE_KERNEL); + + +template +void right_apply_to_dense(std::shared_ptr exec, + const matrix::Diagonal *a, + const matrix::Dense *b, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE( + GKO_DECLARE_DIAGONAL_RIGHT_APPLY_TO_DENSE_KERNEL); + + +template +void apply_to_csr(std::shared_ptr exec, + const matrix::Diagonal *a, + const matrix::Csr *b, + matrix::Csr *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DIAGONAL_APPLY_TO_CSR_KERNEL); + + +template +void right_apply_to_csr(std::shared_ptr exec, + const matrix::Diagonal *a, + const matrix::Csr *b, + matrix::Csr *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DIAGONAL_RIGHT_APPLY_TO_CSR_KERNEL); + + +template +void convert_to_csr(std::shared_ptr exec, + const matrix::Diagonal *source, + matrix::Csr *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DIAGONAL_CONVERT_TO_CSR_KERNEL); + + +template +void conj_transpose(std::shared_ptr exec, + const matrix::Diagonal *orig, + matrix::Diagonal *trans) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DIAGONAL_CONJ_TRANSPOSE_KERNEL); + + +} // namespace diagonal +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/ell_kernels.dp.cpp b/dpcpp/matrix/ell_kernels.dp.cpp new file mode 100644 index 00000000000..f7aedb88628 --- /dev/null +++ b/dpcpp/matrix/ell_kernels.dp.cpp @@ -0,0 +1,142 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/ell_kernels.hpp" + + +#include + + +#include +#include +#include +#include + + +#include "dpcpp/components/format_conversion.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The ELL matrix format namespace. + * + * @ingroup ell + */ +namespace ell { + + +template +void spmv(std::shared_ptr exec, + const matrix::Ell *a, + const matrix::Dense *b, matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_ELL_SPMV_KERNEL); + + +template +void advanced_spmv(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Ell *a, + const matrix::Dense *b, + const matrix::Dense *beta, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ELL_ADVANCED_SPMV_KERNEL); + + +template +void convert_to_dense(std::shared_ptr exec, + const matrix::Ell *source, + matrix::Dense *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ELL_CONVERT_TO_DENSE_KERNEL); + + +template +void convert_to_csr(std::shared_ptr exec, + const matrix::Ell *source, + matrix::Csr *result) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ELL_CONVERT_TO_CSR_KERNEL); + +template +void count_nonzeros(std::shared_ptr exec, + const matrix::Ell *source, + size_type *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ELL_COUNT_NONZEROS_KERNEL); + + +template +void calculate_nonzeros_per_row(std::shared_ptr exec, + const matrix::Ell *source, + Array *result) GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ELL_CALCULATE_NONZEROS_PER_ROW_KERNEL); + + +template +void extract_diagonal(std::shared_ptr exec, + const matrix::Ell *orig, + matrix::Diagonal *diag) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ELL_EXTRACT_DIAGONAL_KERNEL); + + +} // namespace ell +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/hybrid_kernels.dp.cpp b/dpcpp/matrix/hybrid_kernels.dp.cpp new file mode 100644 index 00000000000..e1185553e92 --- /dev/null +++ b/dpcpp/matrix/hybrid_kernels.dp.cpp @@ -0,0 +1,99 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/hybrid_kernels.hpp" + + +#include + + +#include +#include +#include +#include + + +#include "core/matrix/ell_kernels.hpp" +#include "dpcpp/components/format_conversion.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Hybrid matrix format namespace. + * + * @ingroup hybrid + */ +namespace hybrid { + + +template +void convert_to_dense(std::shared_ptr exec, + const matrix::Hybrid *source, + matrix::Dense *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_HYBRID_CONVERT_TO_DENSE_KERNEL); + + +template +void convert_to_csr(std::shared_ptr exec, + const matrix::Hybrid *source, + matrix::Csr *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_HYBRID_CONVERT_TO_CSR_KERNEL); + + +template +void count_nonzeros(std::shared_ptr exec, + const matrix::Hybrid *source, + size_type *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_HYBRID_COUNT_NONZEROS_KERNEL); + + +} // namespace hybrid +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/sellp_kernels.dp.cpp b/dpcpp/matrix/sellp_kernels.dp.cpp new file mode 100644 index 00000000000..20f6adcc2a0 --- /dev/null +++ b/dpcpp/matrix/sellp_kernels.dp.cpp @@ -0,0 +1,125 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/sellp_kernels.hpp" + + +#include + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The SELL-P matrix format namespace. + * + * @ingroup sellp + */ +namespace sellp { + + +template +void spmv(std::shared_ptr exec, + const matrix::Sellp *a, + const matrix::Dense *b, matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SELLP_SPMV_KERNEL); + + +template +void advanced_spmv(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::Sellp *a, + const matrix::Dense *b, + const matrix::Dense *beta, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SELLP_ADVANCED_SPMV_KERNEL); + + +template +void convert_to_dense(std::shared_ptr exec, + const matrix::Sellp *source, + matrix::Dense *result) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SELLP_CONVERT_TO_DENSE_KERNEL); + + +template +void convert_to_csr(std::shared_ptr exec, + const matrix::Sellp *source, + matrix::Csr *result) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SELLP_CONVERT_TO_CSR_KERNEL); + + +template +void count_nonzeros(std::shared_ptr exec, + const matrix::Sellp *source, + size_type *result) GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SELLP_COUNT_NONZEROS_KERNEL); + + +template +void extract_diagonal(std::shared_ptr exec, + const matrix::Sellp *orig, + matrix::Diagonal *diag) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SELLP_EXTRACT_DIAGONAL_KERNEL); + + +} // namespace sellp +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/sparsity_csr_kernels.dp.cpp b/dpcpp/matrix/sparsity_csr_kernels.dp.cpp new file mode 100644 index 00000000000..0093919ccbd --- /dev/null +++ b/dpcpp/matrix/sparsity_csr_kernels.dp.cpp @@ -0,0 +1,175 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/sparsity_csr_kernels.hpp" + + +#include +#include +#include + + +#include + + +#include +#include +#include + + +#include "core/base/iterator_factory.hpp" +#include "dpcpp/components/format_conversion.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The SparsityCsr pattern format namespace. + * + * @ingroup sparsity + */ +namespace sparsity_csr { + + +template +void spmv(std::shared_ptr exec, + const matrix::SparsityCsr *a, + const matrix::Dense *b, matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL); + + +template +void advanced_spmv(std::shared_ptr exec, + const matrix::Dense *alpha, + const matrix::SparsityCsr *a, + const matrix::Dense *b, + const matrix::Dense *beta, + matrix::Dense *c) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); + + +template +void count_num_diagonal_elements( + std::shared_ptr exec, + const matrix::SparsityCsr *matrix, + size_type *num_diagonal_elements) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + + +template +void remove_diagonal_elements(std::shared_ptr exec, + const IndexType *row_ptrs, + const IndexType *col_idxs, + matrix::SparsityCsr *matrix) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + + +template +inline void convert_sparsity_to_csc(size_type num_rows, + const IndexType *row_ptrs, + const IndexType *col_idxs, + IndexType *row_idxs, IndexType *col_ptrs) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void transpose_and_transform( + std::shared_ptr exec, + matrix::SparsityCsr *trans, + const matrix::SparsityCsr *orig) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void transpose(std::shared_ptr exec, + const matrix::SparsityCsr *orig, + matrix::SparsityCsr *trans) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); + + +template +void sort_by_column_index(std::shared_ptr exec, + matrix::SparsityCsr *to_sort) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX); + + +template +void is_sorted_by_column_index( + std::shared_ptr exec, + const matrix::SparsityCsr *to_check, bool *is_sorted) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX); + + +} // namespace sparsity_csr +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/preconditioner/isai_kernels.dp.cpp b/dpcpp/preconditioner/isai_kernels.dp.cpp new file mode 100644 index 00000000000..481979eed3b --- /dev/null +++ b/dpcpp/preconditioner/isai_kernels.dp.cpp @@ -0,0 +1,129 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/preconditioner/isai_kernels.hpp" + + +#include +#include + + +#include + + +#include +#include +#include +#include + + +#include "core/components/prefix_sum.hpp" +#include "core/matrix/csr_builder.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Isai preconditioner namespace. + * + * @ingroup isai + */ +namespace isai { + + +template +void forall_matching(const IndexType *fst, IndexType fst_size, + const IndexType *snd, IndexType snd_size, Callback cb) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void generic_generate(std::shared_ptr exec, + const matrix::Csr *mtx, + matrix::Csr *inverse_mtx, + IndexType *excess_rhs_ptrs, IndexType *excess_nz_ptrs, + Callable trs_solve) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void generate_tri_inverse(std::shared_ptr exec, + const matrix::Csr *mtx, + matrix::Csr *inverse_mtx, + IndexType *excess_rhs_ptrs, IndexType *excess_nz_ptrs, + bool lower) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ISAI_GENERATE_TRI_INVERSE_KERNEL); + + +template +void generate_excess_system(std::shared_ptr, + const matrix::Csr *input, + const matrix::Csr *inverse, + const IndexType *excess_rhs_ptrs, + const IndexType *excess_nz_ptrs, + matrix::Csr *excess_system, + matrix::Dense *excess_rhs) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ISAI_GENERATE_EXCESS_SYSTEM_KERNEL); + + +template +void scatter_excess_solution(std::shared_ptr, + const IndexType *excess_block_ptrs, + const matrix::Dense *excess_solution, + matrix::Csr *inverse) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_ISAI_SCATTER_EXCESS_SOLUTION_KERNEL); + + +} // namespace isai +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/preconditioner/jacobi_kernels.dp.cpp b/dpcpp/preconditioner/jacobi_kernels.dp.cpp new file mode 100644 index 00000000000..0bb653adc65 --- /dev/null +++ b/dpcpp/preconditioner/jacobi_kernels.dp.cpp @@ -0,0 +1,345 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/preconditioner/jacobi_kernels.hpp" + + +#include +#include +#include +#include +#include + + +#include + + +#include +#include +#include +#include + + +#include "core/base/allocator.hpp" +#include "core/base/extended_float.hpp" +#include "core/preconditioner/jacobi_utils.hpp" +#include "dpcpp/components/matrix_operations.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Jacobi preconditioner namespace. + * + * @ingroup jacobi + */ +namespace jacobi { + + +void initialize_precisions(std::shared_ptr exec, + const Array &source, + Array &precisions) +{ + GKO_NOT_IMPLEMENTED; +} + + +namespace { + + +template +inline bool has_same_nonzero_pattern(const IndexType *prev_row_ptr, + const IndexType *curr_row_ptr, + const IndexType *next_row_ptr) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +size_type find_natural_blocks(const matrix::Csr *mtx, + uint32 max_block_size, IndexType *block_ptrs) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +inline size_type agglomerate_supervariables(uint32 max_block_size, + size_type num_natural_blocks, + IndexType *block_ptrs) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace + + +template +void find_blocks(std::shared_ptr exec, + const matrix::Csr *system_matrix, + uint32 max_block_size, size_type &num_blocks, + Array &block_pointers) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_JACOBI_FIND_BLOCKS_KERNEL); + + +namespace { + + +template +inline void extract_block(const matrix::Csr *mtx, + IndexType block_size, IndexType block_start, + ValueType *block, size_type stride) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +inline IndexType choose_pivot(IndexType block_size, const ValueType *block, + size_type stride) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +inline void swap_rows(IndexType row1, IndexType row2, IndexType block_size, + ValueType *block, size_type stride) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +inline bool apply_gauss_jordan_transform(IndexType row, IndexType col, + IndexType block_size, ValueType *block, + size_type stride) +{ + GKO_NOT_IMPLEMENTED; +} + + +template > +inline void transpose_block(IndexType block_size, const SourceValueType *from, + size_type from_stride, ResultValueType *to, + size_type to_stride, + ValueConverter converter = {}) noexcept +{ + GKO_NOT_IMPLEMENTED; +} + + +template > +inline void conj_transpose_block(IndexType block_size, + const SourceValueType *from, + size_type from_stride, ResultValueType *to, + size_type to_stride, + ValueConverter converter = {}) noexcept +{ + GKO_NOT_IMPLEMENTED; +} + + +template > +inline void permute_and_transpose_block(IndexType block_size, + const IndexType *col_perm, + const SourceValueType *source, + size_type source_stride, + ResultValueType *result, + size_type result_stride, + ValueConverter converter = {}) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +inline bool invert_block(IndexType block_size, IndexType *perm, + ValueType *block, size_type stride) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +inline bool validate_precision_reduction_feasibility( + std::shared_ptr exec, IndexType block_size, + const ValueType *block, size_type stride) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace + + +template +void generate(std::shared_ptr exec, + const matrix::Csr *system_matrix, + size_type num_blocks, uint32 max_block_size, + remove_complex accuracy, + const preconditioner::block_interleaved_storage_scheme + &storage_scheme, + Array> &conditioning, + Array &block_precisions, + const Array &block_pointers, Array &blocks) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_JACOBI_GENERATE_KERNEL); + + +namespace { + + +template < + typename ValueType, typename BlockValueType, + typename ValueConverter = default_converter> +inline void apply_block(size_type block_size, size_type num_rhs, + const BlockValueType *block, size_type stride, + ValueType alpha, const ValueType *b, size_type stride_b, + ValueType beta, ValueType *x, size_type stride_x, + ValueConverter converter = {}) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace + + +template +void apply(std::shared_ptr exec, size_type num_blocks, + uint32 max_block_size, + const preconditioner::block_interleaved_storage_scheme + &storage_scheme, + const Array &block_precisions, + const Array &block_pointers, + const Array &blocks, + const matrix::Dense *alpha, + const matrix::Dense *b, + const matrix::Dense *beta, matrix::Dense *x) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_JACOBI_APPLY_KERNEL); + + +template +void simple_apply( + std::shared_ptr exec, size_type num_blocks, + uint32 max_block_size, + const preconditioner::block_interleaved_storage_scheme + &storage_scheme, + const Array &block_precisions, + const Array &block_pointers, const Array &blocks, + const matrix::Dense *b, matrix::Dense *x) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_JACOBI_SIMPLE_APPLY_KERNEL); + + +template +void transpose_jacobi( + std::shared_ptr exec, size_type num_blocks, + uint32 max_block_size, const Array &block_precisions, + const Array &block_pointers, const Array &blocks, + const preconditioner::block_interleaved_storage_scheme + &storage_scheme, + Array &out_blocks) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_JACOBI_TRANSPOSE_KERNEL); + + +template +void conj_transpose_jacobi( + std::shared_ptr exec, size_type num_blocks, + uint32 max_block_size, const Array &block_precisions, + const Array &block_pointers, const Array &blocks, + const preconditioner::block_interleaved_storage_scheme + &storage_scheme, + Array &out_blocks) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_JACOBI_CONJ_TRANSPOSE_KERNEL); + + +template +void convert_to_dense( + std::shared_ptr exec, size_type num_blocks, + const Array &block_precisions, + const Array &block_pointers, const Array &blocks, + const preconditioner::block_interleaved_storage_scheme + &storage_scheme, + ValueType *result_values, size_type result_stride) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_JACOBI_CONVERT_TO_DENSE_KERNEL); + + +} // namespace jacobi +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/bicg_kernels.dp.cpp b/dpcpp/solver/bicg_kernels.dp.cpp new file mode 100644 index 00000000000..11a5471002a --- /dev/null +++ b/dpcpp/solver/bicg_kernels.dp.cpp @@ -0,0 +1,105 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/bicg_kernels.hpp" + + +#include + + +#include +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The BICG solver namespace. + * + * @ingroup bicg + */ +namespace bicg { + + +template +void initialize(std::shared_ptr exec, + const matrix::Dense *b, matrix::Dense *r, + matrix::Dense *z, matrix::Dense *p, + matrix::Dense *q, matrix::Dense *prev_rho, + matrix::Dense *rho, matrix::Dense *r2, + matrix::Dense *z2, matrix::Dense *p2, + matrix::Dense *q2, + Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICG_INITIALIZE_KERNEL); + + +template +void step_1(std::shared_ptr exec, + matrix::Dense *p, const matrix::Dense *z, + matrix::Dense *p2, const matrix::Dense *z2, + const matrix::Dense *rho, + const matrix::Dense *prev_rho, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICG_STEP_1_KERNEL); + + +template +void step_2(std::shared_ptr exec, + matrix::Dense *x, matrix::Dense *r, + matrix::Dense *r2, const matrix::Dense *p, + const matrix::Dense *q, + const matrix::Dense *q2, + const matrix::Dense *beta, + const matrix::Dense *rho, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICG_STEP_2_KERNEL); + + +} // namespace bicg +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/bicgstab_kernels.dp.cpp b/dpcpp/solver/bicgstab_kernels.dp.cpp new file mode 100644 index 00000000000..af4d42289cd --- /dev/null +++ b/dpcpp/solver/bicgstab_kernels.dp.cpp @@ -0,0 +1,137 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/bicgstab_kernels.hpp" + + +#include + + +#include + + +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The BICGSTAB solver namespace. + * + * @ingroup bicgstab + */ +namespace bicgstab { + + +template +void initialize(std::shared_ptr exec, + const matrix::Dense *b, matrix::Dense *r, + matrix::Dense *rr, matrix::Dense *y, + matrix::Dense *s, matrix::Dense *t, + matrix::Dense *z, matrix::Dense *v, + matrix::Dense *p, matrix::Dense *prev_rho, + matrix::Dense *rho, matrix::Dense *alpha, + matrix::Dense *beta, matrix::Dense *gamma, + matrix::Dense *omega, + Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICGSTAB_INITIALIZE_KERNEL); + + +template +void step_1(std::shared_ptr exec, + const matrix::Dense *r, matrix::Dense *p, + const matrix::Dense *v, + const matrix::Dense *rho, + const matrix::Dense *prev_rho, + const matrix::Dense *alpha, + const matrix::Dense *omega, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICGSTAB_STEP_1_KERNEL); + + +template +void step_2(std::shared_ptr exec, + const matrix::Dense *r, matrix::Dense *s, + const matrix::Dense *v, + const matrix::Dense *rho, + matrix::Dense *alpha, + const matrix::Dense *beta, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICGSTAB_STEP_2_KERNEL); + + +template +void step_3( + std::shared_ptr exec, matrix::Dense *x, + matrix::Dense *r, const matrix::Dense *s, + const matrix::Dense *t, const matrix::Dense *y, + const matrix::Dense *z, const matrix::Dense *alpha, + const matrix::Dense *beta, const matrix::Dense *gamma, + matrix::Dense *omega, const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICGSTAB_STEP_3_KERNEL); + + +template +void finalize(std::shared_ptr exec, + matrix::Dense *x, const matrix::Dense *y, + const matrix::Dense *alpha, + Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BICGSTAB_FINALIZE_KERNEL); + + +} // namespace bicgstab +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/cg_kernels.dp.cpp b/dpcpp/solver/cg_kernels.dp.cpp new file mode 100644 index 00000000000..9cb2e0f30da --- /dev/null +++ b/dpcpp/solver/cg_kernels.dp.cpp @@ -0,0 +1,101 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/cg_kernels.hpp" + + +#include + + +#include +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The CG solver namespace. + * + * @ingroup cg + */ +namespace cg { + + +template +void initialize(std::shared_ptr exec, + const matrix::Dense *b, matrix::Dense *r, + matrix::Dense *z, matrix::Dense *p, + matrix::Dense *q, matrix::Dense *prev_rho, + matrix::Dense *rho, + Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CG_INITIALIZE_KERNEL); + + +template +void step_1(std::shared_ptr exec, + matrix::Dense *p, const matrix::Dense *z, + const matrix::Dense *rho, + const matrix::Dense *prev_rho, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CG_STEP_1_KERNEL); + + +template +void step_2(std::shared_ptr exec, + matrix::Dense *x, matrix::Dense *r, + const matrix::Dense *p, + const matrix::Dense *q, + const matrix::Dense *beta, + const matrix::Dense *rho, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CG_STEP_2_KERNEL); + + +} // namespace cg +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/cgs_kernels.dp.cpp b/dpcpp/solver/cgs_kernels.dp.cpp new file mode 100644 index 00000000000..04d81b682c5 --- /dev/null +++ b/dpcpp/solver/cgs_kernels.dp.cpp @@ -0,0 +1,119 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/cgs_kernels.hpp" + + +#include + + +#include +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The CGS solver namespace. + * + * @ingroup cgs + */ +namespace cgs { + + +template +void initialize(std::shared_ptr exec, + const matrix::Dense *b, matrix::Dense *r, + matrix::Dense *r_tld, matrix::Dense *p, + matrix::Dense *q, matrix::Dense *u, + matrix::Dense *u_hat, + matrix::Dense *v_hat, matrix::Dense *t, + matrix::Dense *alpha, matrix::Dense *beta, + matrix::Dense *gamma, + matrix::Dense *prev_rho, + matrix::Dense *rho, + Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CGS_INITIALIZE_KERNEL); + + +template +void step_1(std::shared_ptr exec, + const matrix::Dense *r, matrix::Dense *u, + matrix::Dense *p, const matrix::Dense *q, + matrix::Dense *beta, const matrix::Dense *rho, + const matrix::Dense *rho_prev, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CGS_STEP_1_KERNEL); + + +template +void step_2(std::shared_ptr exec, + const matrix::Dense *u, + const matrix::Dense *v_hat, matrix::Dense *q, + matrix::Dense *t, matrix::Dense *alpha, + const matrix::Dense *rho, + const matrix::Dense *gamma, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CGS_STEP_2_KERNEL); + +template +void step_3(std::shared_ptr exec, + const matrix::Dense *t, + const matrix::Dense *u_hat, matrix::Dense *r, + matrix::Dense *x, const matrix::Dense *alpha, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_CGS_STEP_3_KERNEL); + + +} // namespace cgs +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/fcg_kernels.dp.cpp b/dpcpp/solver/fcg_kernels.dp.cpp new file mode 100644 index 00000000000..79422ed1960 --- /dev/null +++ b/dpcpp/solver/fcg_kernels.dp.cpp @@ -0,0 +1,101 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/fcg_kernels.hpp" + + +#include + + +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The FCG solver namespace. + * + * @ingroup fcg + */ +namespace fcg { + + +template +void initialize(std::shared_ptr exec, + const matrix::Dense *b, matrix::Dense *r, + matrix::Dense *z, matrix::Dense *p, + matrix::Dense *q, matrix::Dense *t, + matrix::Dense *prev_rho, + matrix::Dense *rho, matrix::Dense *rho_t, + Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_FCG_INITIALIZE_KERNEL); + + +template +void step_1(std::shared_ptr exec, + matrix::Dense *p, const matrix::Dense *z, + const matrix::Dense *rho_t, + const matrix::Dense *prev_rho, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_FCG_STEP_1_KERNEL); + + +template +void step_2(std::shared_ptr exec, + matrix::Dense *x, matrix::Dense *r, + matrix::Dense *t, const matrix::Dense *p, + const matrix::Dense *q, + const matrix::Dense *beta, + const matrix::Dense *rho, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_FCG_STEP_2_KERNEL); + + +} // namespace fcg +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/gmres_kernels.dp.cpp b/dpcpp/solver/gmres_kernels.dp.cpp new file mode 100644 index 00000000000..c8924b0b67f --- /dev/null +++ b/dpcpp/solver/gmres_kernels.dp.cpp @@ -0,0 +1,186 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/gmres_kernels.hpp" + + +#include + + +#include +#include +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The GMRES solver namespace. + * + * @ingroup gmres + */ +namespace gmres { + + +namespace { + + +template +void finish_arnoldi(size_type num_rows, matrix::Dense *krylov_bases, + matrix::Dense *hessenberg_iter, size_type iter, + const stopping_status *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void calculate_sin_and_cos(matrix::Dense *givens_sin, + matrix::Dense *givens_cos, + matrix::Dense *hessenberg_iter, + size_type iter, const size_type rhs) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void givens_rotation(matrix::Dense *givens_sin, + matrix::Dense *givens_cos, + matrix::Dense *hessenberg_iter, size_type iter, + const stopping_status *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void calculate_next_residual_norm( + matrix::Dense *givens_sin, matrix::Dense *givens_cos, + matrix::Dense> *residual_norm, + matrix::Dense *residual_norm_collection, size_type iter, + const stopping_status *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void solve_upper_triangular( + const matrix::Dense *residual_norm_collection, + const matrix::Dense *hessenberg, matrix::Dense *y, + const size_type *final_iter_nums) +{ + GKO_NOT_IMPLEMENTED; +} + + +template +void calculate_qy(const matrix::Dense *krylov_bases, + const matrix::Dense *y, + matrix::Dense *before_preconditioner, + const size_type *final_iter_nums) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace + + +template +void initialize_1(std::shared_ptr exec, + const matrix::Dense *b, + matrix::Dense *residual, + matrix::Dense *givens_sin, + matrix::Dense *givens_cos, + Array *stop_status, size_type krylov_dim) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GMRES_INITIALIZE_1_KERNEL); + + +template +void initialize_2(std::shared_ptr exec, + const matrix::Dense *residual, + matrix::Dense> *residual_norm, + matrix::Dense *residual_norm_collection, + matrix::Dense *krylov_bases, + Array *final_iter_nums, size_type krylov_dim) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GMRES_INITIALIZE_2_KERNEL); + + +template +void step_1(std::shared_ptr exec, size_type num_rows, + matrix::Dense *givens_sin, + matrix::Dense *givens_cos, + matrix::Dense> *residual_norm, + matrix::Dense *residual_norm_collection, + matrix::Dense *krylov_bases, + matrix::Dense *hessenberg_iter, size_type iter, + Array *final_iter_nums, + const Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GMRES_STEP_1_KERNEL); + + +template +void step_2(std::shared_ptr exec, + const matrix::Dense *residual_norm_collection, + const matrix::Dense *krylov_bases, + const matrix::Dense *hessenberg, + matrix::Dense *y, + matrix::Dense *before_preconditioner, + const Array *final_iter_nums) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GMRES_STEP_2_KERNEL); + + +} // namespace gmres +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/ir_kernels.dp.cpp b/dpcpp/solver/ir_kernels.dp.cpp new file mode 100644 index 00000000000..bd3c6a98bc8 --- /dev/null +++ b/dpcpp/solver/ir_kernels.dp.cpp @@ -0,0 +1,60 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/ir_kernels.hpp" + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The IR solver namespace. + * + * @ingroup ir + */ +namespace ir { + + +void initialize(std::shared_ptr exec, + Array *stop_status) +{ + GKO_NOT_IMPLEMENTED; +} + + +} // namespace ir +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/lower_trs_kernels.dp.cpp b/dpcpp/solver/lower_trs_kernels.dp.cpp new file mode 100644 index 00000000000..6f0546ed6a4 --- /dev/null +++ b/dpcpp/solver/lower_trs_kernels.dp.cpp @@ -0,0 +1,112 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/lower_trs_kernels.hpp" + + +#include + + +#include + + +#include +#include +#include +#include +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The LOWER_TRS solver namespace. + * + * @ingroup lower_trs + */ +namespace lower_trs { + + +void should_perform_transpose(std::shared_ptr exec, + bool &do_transpose) +{ + GKO_NOT_IMPLEMENTED; +} + + +void init_struct(std::shared_ptr exec, + std::shared_ptr &solve_struct) +{ + // This init kernel is here to allow initialization of the solve struct for + // a more sophisticated implementation as for other executors. +} + + +template +void generate(std::shared_ptr exec, + const matrix::Csr *matrix, + solver::SolveStruct *solve_struct, const gko::size_type num_rhs) +{ + // This generate kernel is here to allow for a more sophisticated + // implementation as for other executors. This kernel would perform the + // "analysis" phase for the triangular matrix. +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_LOWER_TRS_GENERATE_KERNEL); + + +/** + * The parameters trans_x and trans_b are used only in the CUDA executor for + * versions <=9.1 due to a limitation in the cssrsm_solve algorithm + */ +template +void solve(std::shared_ptr exec, + const matrix::Csr *matrix, + const solver::SolveStruct *solve_struct, + matrix::Dense *trans_b, matrix::Dense *trans_x, + const matrix::Dense *b, matrix::Dense *x) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_LOWER_TRS_SOLVE_KERNEL); + + +} // namespace lower_trs +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/solver/upper_trs_kernels.dp.cpp b/dpcpp/solver/upper_trs_kernels.dp.cpp new file mode 100644 index 00000000000..67efc9896f6 --- /dev/null +++ b/dpcpp/solver/upper_trs_kernels.dp.cpp @@ -0,0 +1,112 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/solver/upper_trs_kernels.hpp" + + +#include + + +#include + + +#include +#include +#include +#include +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The UPPER_TRS solver namespace. + * + * @ingroup upper_trs + */ +namespace upper_trs { + + +void should_perform_transpose(std::shared_ptr exec, + bool &do_transpose) +{ + do_transpose = false; +} + + +void init_struct(std::shared_ptr exec, + std::shared_ptr &solve_struct) +{ + // This init kernel is here to allow initialization of the solve struct for + // a more sophisticated implementation as for other executors. +} + + +template +void generate(std::shared_ptr exec, + const matrix::Csr *matrix, + solver::SolveStruct *solve_struct, const gko::size_type num_rhs) +{ + // This generate kernel is here to allow for a more sophisticated + // implementation as for other executors. This kernel would perform the + // "analysis" phase for the triangular matrix. +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_UPPER_TRS_GENERATE_KERNEL); + + +/** + * The parameters trans_x and trans_b are used only in the CUDA executor for + * versions <=9.1 due to a limitation in the cssrsm_solve algorithm + */ +template +void solve(std::shared_ptr exec, + const matrix::Csr *matrix, + const solver::SolveStruct *solve_struct, + matrix::Dense *trans_b, matrix::Dense *trans_x, + const matrix::Dense *b, matrix::Dense *x) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_UPPER_TRS_SOLVE_KERNEL); + + +} // namespace upper_trs +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/stop/criterion_kernels.dp.cpp b/dpcpp/stop/criterion_kernels.dp.cpp new file mode 100644 index 00000000000..dce2c8ae84c --- /dev/null +++ b/dpcpp/stop/criterion_kernels.dp.cpp @@ -0,0 +1,71 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/stop/criterion_kernels.hpp" + + +#include + + +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Setting of all statuses namespace. + * @ref set_status + * @ingroup set_all_statuses + */ +namespace set_all_statuses { + + +void set_all_statuses(std::shared_ptr exec, + uint8 stoppingId, bool setFinalized, + Array *stop_status) +{ + auto size = stop_status->get_num_elems(); + stopping_status *__restrict__ stop_status_ptr = stop_status->get_data(); + exec->get_queue()->submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<1>{size}, [=](sycl::id<1> idx_id) { + const int idx = idx_id[0]; + stop_status_ptr[idx].stop(stoppingId, setFinalized); + }); + }); +} + + +} // namespace set_all_statuses +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/stop/residual_norm_kernels.dp.cpp b/dpcpp/stop/residual_norm_kernels.dp.cpp new file mode 100644 index 00000000000..f0f205d438d --- /dev/null +++ b/dpcpp/stop/residual_norm_kernels.dp.cpp @@ -0,0 +1,77 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/stop/residual_norm_kernels.hpp" + + +#include + + +#include +#include +#include + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Residual norm stopping criterion namespace. + * @ref resnorm + * @ingroup resnorm + */ +namespace residual_norm { + + +constexpr int default_group_size = 512; + + +template +void residual_norm(std::shared_ptr exec, + const matrix::Dense *tau, + const matrix::Dense *orig_tau, + ValueType rel_residual_goal, uint8 stoppingId, + bool setFinalized, Array *stop_status, + Array *device_storage, bool *all_converged, + bool *one_changed) +{ + GKO_NOT_IMPLEMENTED; +} + +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_TYPE( + GKO_DECLARE_RESIDUAL_NORM_KERNEL); + + +} // namespace residual_norm +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/test/CMakeLists.txt b/dpcpp/test/CMakeLists.txt new file mode 100644 index 00000000000..8e34766a1c5 --- /dev/null +++ b/dpcpp/test/CMakeLists.txt @@ -0,0 +1,5 @@ +include(${PROJECT_SOURCE_DIR}/cmake/create_test.cmake) + +add_subdirectory(base) +add_subdirectory(components) +add_subdirectory(stop) diff --git a/dpcpp/test/base/CMakeLists.txt b/dpcpp/test/base/CMakeLists.txt new file mode 100644 index 00000000000..adb2f2505b5 --- /dev/null +++ b/dpcpp/test/base/CMakeLists.txt @@ -0,0 +1 @@ +ginkgo_create_dpcpp_test(executor) diff --git a/dpcpp/test/base/executor.dp.cpp b/dpcpp/test/base/executor.dp.cpp new file mode 100644 index 00000000000..49cf10735a7 --- /dev/null +++ b/dpcpp/test/base/executor.dp.cpp @@ -0,0 +1,303 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include + + +#include +#include +#include + + +#include + + +#include + + +#include +#include +#include + + +namespace { + + +class DpcppExecutor : public ::testing::Test { +protected: + DpcppExecutor() + : omp(gko::OmpExecutor::create()), dpcpp(nullptr), dpcpp2(nullptr) + {} + + void SetUp() + { + ASSERT_GT(gko::DpcppExecutor::get_num_devices("cpu"), 0); + dpcpp = gko::DpcppExecutor::create(0, omp, "cpu"); + if (gko::DpcppExecutor::get_num_devices("gpu") > 0) { + dpcpp2 = gko::DpcppExecutor::create(0, omp, "gpu"); + } + } + + void TearDown() + { + if (dpcpp != nullptr) { + // ensure that previous calls finished and didn't throw an error + ASSERT_NO_THROW(dpcpp->synchronize()); + } + } + + std::shared_ptr omp; + std::shared_ptr dpcpp; + std::shared_ptr dpcpp2; +}; + + +TEST_F(DpcppExecutor, CanInstantiateTwoExecutorsOnOneDevice) +{ + auto dpcpp = gko::DpcppExecutor::create(0, omp, "all"); + auto dpcpp2 = gko::DpcppExecutor::create(0, omp, "all"); + + // We want automatic deinitialization to not create any error +} + + +TEST_F(DpcppExecutor, KnowsNumberOfDevicesOfTypeAll) +{ + auto count = sycl::device::get_devices(sycl::info::device_type::all).size(); + + auto num_devices = gko::DpcppExecutor::get_num_devices("all"); + + ASSERT_EQ(count, num_devices); +} + + +TEST_F(DpcppExecutor, KnowsNumberOfDevicesOfTypeCPU) +{ + auto count = sycl::device::get_devices(sycl::info::device_type::cpu).size(); + + auto num_devices = gko::DpcppExecutor::get_num_devices("cpu"); + + ASSERT_EQ(count, num_devices); +} + + +TEST_F(DpcppExecutor, KnowsNumberOfDevicesOfTypeGPU) +{ + auto count = sycl::device::get_devices(sycl::info::device_type::gpu).size(); + + auto num_devices = gko::DpcppExecutor::get_num_devices("gpu"); + + ASSERT_EQ(count, num_devices); +} + + +TEST_F(DpcppExecutor, KnowsNumberOfDevicesOfTypeAccelerator) +{ + auto count = + sycl::device::get_devices(sycl::info::device_type::accelerator).size(); + + auto num_devices = gko::DpcppExecutor::get_num_devices("accelerator"); + + ASSERT_EQ(count, num_devices); +} + + +TEST_F(DpcppExecutor, AllocatesAndFreesMemoryOnCPU) +{ + int *ptr = nullptr; + + ASSERT_NO_THROW(ptr = dpcpp->alloc(2)); + ASSERT_NO_THROW(dpcpp->free(ptr)); +} + + +TEST_F(DpcppExecutor, AllocatesAndFreesMemoryOnGPU) +{ + if (!dpcpp2) { + GTEST_SKIP() << "No DPC++ compatible GPU."; + } + int *ptr = nullptr; + + ASSERT_NO_THROW(ptr = dpcpp2->alloc(2)); + ASSERT_NO_THROW(dpcpp2->free(ptr)); +} + + +TEST_F(DpcppExecutor, FailsWhenOverallocating) +{ + const gko::size_type num_elems = 1ll << 50; // 4PB of integers + int *ptr = nullptr; + + ASSERT_THROW( + { + ptr = dpcpp->alloc(num_elems); + dpcpp->synchronize(); + }, + gko::AllocationError); + + dpcpp->free(ptr); +} + + +void check_data(int *data, bool *result) +{ + *result = false; + if (data[0] == 3 && data[1] == 8) { + *result = true; + } +} + +TEST_F(DpcppExecutor, CopiesDataToCPU) +{ + int orig[] = {3, 8}; + auto *copy = dpcpp->alloc(2); + gko::Array is_set(omp, 1); + + dpcpp->copy_from(omp.get(), 2, orig, copy); + + is_set.set_executor(dpcpp); + ASSERT_NO_THROW(dpcpp->synchronize()); + ASSERT_NO_THROW(dpcpp->get_queue()->submit([&](sycl::handler &cgh) { + auto *is_set_ptr = is_set.get_data(); + cgh.single_task([=]() { check_data(copy, is_set_ptr); }); + })); + is_set.set_executor(omp); + ASSERT_EQ(*is_set.get_data(), true); + ASSERT_NO_THROW(dpcpp->synchronize()); + dpcpp->free(copy); +} + + +TEST_F(DpcppExecutor, CopiesDataToGPU) +{ + if (!dpcpp2) { + GTEST_SKIP() << "No DPC++ compatible GPU."; + } + int orig[] = {3, 8}; + auto *copy = dpcpp2->alloc(2); + gko::Array is_set(omp, 1); + + dpcpp2->copy_from(omp.get(), 2, orig, copy); + + is_set.set_executor(dpcpp2); + ASSERT_NO_THROW(dpcpp2->get_queue()->submit([&](sycl::handler &cgh) { + auto *is_set_ptr = is_set.get_data(); + cgh.single_task([=]() { check_data(copy, is_set_ptr); }); + })); + is_set.set_executor(omp); + ASSERT_EQ(*is_set.get_data(), true); + ASSERT_NO_THROW(dpcpp2->synchronize()); + dpcpp2->free(copy); +} + + +void init_data(int *data) +{ + data[0] = 3; + data[1] = 8; +} + +TEST_F(DpcppExecutor, CopiesDataFromCPU) +{ + int copy[2]; + auto orig = dpcpp->alloc(2); + dpcpp->get_queue()->submit([&](sycl::handler &cgh) { + cgh.single_task([=]() { init_data(orig); }); + }); + + omp->copy_from(dpcpp.get(), 2, orig, copy); + + EXPECT_EQ(3, copy[0]); + ASSERT_EQ(8, copy[1]); + dpcpp->free(orig); +} + + +TEST_F(DpcppExecutor, CopiesDataFromGPU) +{ + if (!dpcpp2) { + GTEST_SKIP() << "No DPC++ compatible GPU."; + } + int copy[2]; + auto orig = dpcpp2->alloc(2); + dpcpp2->get_queue()->submit([&](sycl::handler &cgh) { + cgh.single_task([=]() { init_data(orig); }); + }); + + omp->copy_from(dpcpp2.get(), 2, orig, copy); + + EXPECT_EQ(3, copy[0]); + ASSERT_EQ(8, copy[1]); + dpcpp2->free(orig); +} + + +TEST_F(DpcppExecutor, CopiesDataFromDpcppToDpcpp) +{ + if (!dpcpp2) { + GTEST_SKIP() << "No DPC++ compatible GPU."; + } + int copy[2]; + gko::Array is_set(omp, 1); + auto orig = dpcpp->alloc(2); + dpcpp->get_queue()->submit([&](sycl::handler &cgh) { + cgh.single_task([=]() { init_data(orig); }); + }); + + auto copy_dpcpp2 = dpcpp2->alloc(2); + dpcpp2->copy_from(dpcpp.get(), 2, orig, copy_dpcpp2); + // Check that the data is really on GPU + is_set.set_executor(dpcpp2); + ASSERT_NO_THROW(dpcpp2->get_queue()->submit([&](sycl::handler &cgh) { + auto *is_set_ptr = is_set.get_data(); + cgh.single_task([=]() { check_data(copy_dpcpp2, is_set_ptr); }); + })); + is_set.set_executor(omp); + ASSERT_EQ(*is_set.get_data(), true); + + // Put the results on OpenMP and run CPU side assertions + omp->copy_from(dpcpp2.get(), 2, copy_dpcpp2, copy); + EXPECT_EQ(3, copy[0]); + ASSERT_EQ(8, copy[1]); + dpcpp2->free(copy_dpcpp2); + dpcpp->free(orig); +} + + +TEST_F(DpcppExecutor, Synchronizes) +{ + // Todo design a proper unit test once we support streams + ASSERT_NO_THROW(dpcpp->synchronize()); +} + + +} // namespace diff --git a/dpcpp/test/components/CMakeLists.txt b/dpcpp/test/components/CMakeLists.txt new file mode 100644 index 00000000000..f6ca175d105 --- /dev/null +++ b/dpcpp/test/components/CMakeLists.txt @@ -0,0 +1,3 @@ +ginkgo_create_test(absolute_array) +ginkgo_create_test(fill_array) +ginkgo_create_test(precision_conversion) diff --git a/dpcpp/test/components/absolute_array.cpp b/dpcpp/test/components/absolute_array.cpp new file mode 100644 index 00000000000..7dcd1547432 --- /dev/null +++ b/dpcpp/test/components/absolute_array.cpp @@ -0,0 +1,132 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/components/absolute_array.hpp" + + +#include +#include +#include + + +#include + + +#include + + +#include "core/test/utils.hpp" + + +namespace { + + +class AbsoluteArray : public ::testing::Test { +protected: + using value_type = double; + using complex_type = std::complex; + AbsoluteArray() + : ref(gko::ReferenceExecutor::create()), + exec(gko::DpcppExecutor::create(0, ref, "all")), + total_size(6344), + vals(ref, total_size), + dvals(exec, total_size), + complex_vals(ref, total_size), + dcomplex_vals(exec, total_size) + { + std::fill_n(vals.get_data(), total_size, -1234.0); + dvals = vals; + std::fill_n(complex_vals.get_data(), total_size, complex_type{3, 4}); + dcomplex_vals = complex_vals; + } + + std::shared_ptr ref; + std::shared_ptr exec; + gko::size_type total_size; + gko::Array vals; + gko::Array dvals; + gko::Array complex_vals; + gko::Array dcomplex_vals; +}; + + +TEST_F(AbsoluteArray, InplaceEqualsReference) +{ + gko::kernels::dpcpp::components::inplace_absolute_array( + exec, dvals.get_data(), total_size); + gko::kernels::reference::components::inplace_absolute_array( + ref, vals.get_data(), total_size); + + GKO_ASSERT_ARRAY_EQ(vals, dvals); +} + + +TEST_F(AbsoluteArray, InplaceComplexEqualsReference) +{ + gko::kernels::dpcpp::components::inplace_absolute_array( + exec, dcomplex_vals.get_data(), total_size); + gko::kernels::reference::components::inplace_absolute_array( + ref, complex_vals.get_data(), total_size); + + GKO_ASSERT_ARRAY_EQ(complex_vals, dcomplex_vals); +} + + +TEST_F(AbsoluteArray, OutplaceEqualsReference) +{ + gko::Array abs_vals(ref, total_size); + gko::Array dabs_vals(exec, total_size); + + gko::kernels::dpcpp::components::outplace_absolute_array( + exec, dvals.get_const_data(), total_size, dabs_vals.get_data()); + gko::kernels::reference::components::outplace_absolute_array( + ref, vals.get_const_data(), total_size, abs_vals.get_data()); + + GKO_ASSERT_ARRAY_EQ(abs_vals, dabs_vals); +} + + +TEST_F(AbsoluteArray, OutplaceComplexEqualsReference) +{ + gko::Array abs_vals(ref, total_size); + gko::Array dabs_vals(exec, total_size); + + gko::kernels::dpcpp::components::outplace_absolute_array( + exec, dcomplex_vals.get_const_data(), total_size, dabs_vals.get_data()); + gko::kernels::reference::components::outplace_absolute_array( + ref, complex_vals.get_const_data(), total_size, abs_vals.get_data()); + + GKO_ASSERT_ARRAY_EQ(abs_vals, dabs_vals); +} + + +} // namespace diff --git a/dpcpp/test/components/fill_array.cpp b/dpcpp/test/components/fill_array.cpp new file mode 100644 index 00000000000..3122a3b9625 --- /dev/null +++ b/dpcpp/test/components/fill_array.cpp @@ -0,0 +1,86 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/components/fill_array.hpp" + + +#include +#include +#include + + +#include + + +#include + + +#include "core/test/utils.hpp" + + +namespace { + + +template +class FillArray : public ::testing::Test { +protected: + using value_type = T; + FillArray() + : ref(gko::ReferenceExecutor::create()), + exec(gko::DpcppExecutor::create(0, ref, "all")), + total_size(63531), + vals(ref, total_size), + dvals(exec, total_size) + { + std::fill_n(vals.get_data(), total_size, T(1523)); + } + + std::shared_ptr ref; + std::shared_ptr exec; + gko::size_type total_size; + gko::Array vals; + gko::Array dvals; +}; + +TYPED_TEST_SUITE(FillArray, gko::test::ValueAndIndexTypes); + + +TYPED_TEST(FillArray, EqualsReference) +{ + using T = typename TestFixture::value_type; + gko::kernels::dpcpp::components::fill_array( + this->exec, this->dvals.get_data(), this->total_size, T(1523)); + GKO_ASSERT_ARRAY_EQ(this->vals, this->dvals); +} + + +} // namespace diff --git a/dpcpp/test/components/precision_conversion.cpp b/dpcpp/test/components/precision_conversion.cpp new file mode 100644 index 00000000000..36a589b74ed --- /dev/null +++ b/dpcpp/test/components/precision_conversion.cpp @@ -0,0 +1,173 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include +#include +#include +#include +#include + + +#include + + +#include + + +#include "core/test/utils.hpp" + + +namespace { + + +class PrecisionConversion : public ::testing::Test { +protected: + PrecisionConversion() + : ref(gko::ReferenceExecutor::create()), + exec(gko::DpcppExecutor::create(0, ref, "all")), + rand(293), + total_size(42793), + vals(ref, total_size), + cvals(ref, total_size), + vals2(ref, 1), + expected_float(ref, 1), + expected_double(ref, 1), + dvals(exec), + dcvals(exec), + dvals2(exec) + { + auto maxval = 1e10f; + std::uniform_real_distribution dist(-maxval, maxval); + for (gko::size_type i = 0; i < total_size; ++i) { + vals.get_data()[i] = dist(rand); + cvals.get_data()[i] = {dist(rand), dist(rand)}; + } + dvals = vals; + dcvals = cvals; + gko::uint64 rawdouble{0x4218888000889111ULL}; + gko::uint32 rawfloat{0x50c44400UL}; + gko::uint64 rawrounded{0x4218888000000000ULL}; + std::memcpy(vals2.get_data(), &rawdouble, sizeof(double)); + std::memcpy(expected_float.get_data(), &rawfloat, sizeof(float)); + std::memcpy(expected_double.get_data(), &rawrounded, sizeof(double)); + dvals2 = vals2; + } + + std::shared_ptr ref; + std::shared_ptr exec; + std::default_random_engine rand; + gko::size_type total_size; + gko::Array vals; + gko::Array dvals; + gko::Array vals2; + gko::Array dvals2; + gko::Array expected_float; + gko::Array expected_double; + gko::Array> cvals; + gko::Array> dcvals; +}; + + +TEST_F(PrecisionConversion, ConvertsReal) +{ + gko::Array dtmp; + gko::Array dout; + + dtmp = dvals; + dout = dtmp; + + GKO_ASSERT_ARRAY_EQ(dvals, dout); +} + + +TEST_F(PrecisionConversion, ConvertsRealViaRef) +{ + gko::Array tmp{ref}; + gko::Array dout; + + tmp = dvals; + dout = tmp; + + GKO_ASSERT_ARRAY_EQ(dvals, dout); +} + + +TEST_F(PrecisionConversion, ConvertsComplex) +{ + gko::Array> dtmp; + gko::Array> dout; + + dtmp = dcvals; + dout = dtmp; + + GKO_ASSERT_ARRAY_EQ(dcvals, dout); +} + + +TEST_F(PrecisionConversion, ConversionRounds) +{ + gko::Array dtmp; + gko::Array dout; + + dtmp = dvals2; + dout = dtmp; + + GKO_ASSERT_ARRAY_EQ(dtmp, expected_float); + GKO_ASSERT_ARRAY_EQ(dout, expected_double); +} + + +TEST_F(PrecisionConversion, ConvertsRealFromRef) +{ + gko::Array dtmp; + gko::Array dout; + + dtmp = vals; + dout = dtmp; + + GKO_ASSERT_ARRAY_EQ(dvals, dout); +} + + +TEST_F(PrecisionConversion, ConvertsComplexFromRef) +{ + gko::Array> dtmp; + gko::Array> dout; + + dtmp = cvals; + dout = dtmp; + + GKO_ASSERT_ARRAY_EQ(dcvals, dout); +} + + +} // namespace diff --git a/dpcpp/test/stop/CMakeLists.txt b/dpcpp/test/stop/CMakeLists.txt new file mode 100644 index 00000000000..dd96e64d469 --- /dev/null +++ b/dpcpp/test/stop/CMakeLists.txt @@ -0,0 +1 @@ +ginkgo_create_test(criterion_kernels) diff --git a/dpcpp/test/stop/criterion_kernels.cpp b/dpcpp/test/stop/criterion_kernels.cpp new file mode 100644 index 00000000000..09612862566 --- /dev/null +++ b/dpcpp/test/stop/criterion_kernels.cpp @@ -0,0 +1,107 @@ +/************************************************************* +Copyright (c) 2017-2020, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include + + +#include + + +#include + + +namespace { + + +constexpr gko::size_type test_iterations = 10; + + +class Criterion : public ::testing::Test { +protected: + Criterion() + { + ref_ = gko::ReferenceExecutor::create(); + dpcpp_ = gko::DpcppExecutor::create(0, ref_, "all"); + // Actually use an iteration stopping criterion because Criterion is an + // abstract class + factory_ = gko::stop::Iteration::build() + .with_max_iters(test_iterations) + .on(dpcpp_); + } + + std::unique_ptr factory_; + std::shared_ptr ref_; + std::shared_ptr dpcpp_; +}; + + +TEST_F(Criterion, SetsOneStopStatus) +{ + bool one_changed{}; + constexpr gko::uint8 RelativeStoppingId{1}; + auto criterion = factory_->generate(nullptr, nullptr, nullptr); + gko::Array stop_status(ref_, 1); + stop_status.get_data()[0].reset(); + + stop_status.set_executor(dpcpp_); + criterion->update() + .num_iterations(test_iterations) + .check(RelativeStoppingId, true, &stop_status, &one_changed); + stop_status.set_executor(ref_); + + ASSERT_EQ(stop_status.get_data()[0].has_stopped(), true); +} + + +TEST_F(Criterion, SetsMultipleStopStatuses) +{ + bool one_changed{}; + constexpr gko::uint8 RelativeStoppingId{1}; + auto criterion = factory_->generate(nullptr, nullptr, nullptr); + gko::Array stop_status(ref_, 3); + stop_status.get_data()[0].reset(); + stop_status.get_data()[1].reset(); + stop_status.get_data()[2].reset(); + + stop_status.set_executor(dpcpp_); + criterion->update() + .num_iterations(test_iterations) + .check(RelativeStoppingId, true, &stop_status, &one_changed); + stop_status.set_executor(ref_); + + ASSERT_EQ(stop_status.get_data()[0].has_stopped(), true); + ASSERT_EQ(stop_status.get_data()[1].has_stopped(), true); + ASSERT_EQ(stop_status.get_data()[2].has_stopped(), true); +} + + +} // namespace diff --git a/dpcpp/test_dpcpp.dp.cpp b/dpcpp/test_dpcpp.dp.cpp new file mode 100644 index 00000000000..2f63ffdfda2 --- /dev/null +++ b/dpcpp/test_dpcpp.dp.cpp @@ -0,0 +1,10 @@ +#include + +namespace sycl = cl::sycl; + +int main() +{ + // Use the queue property `in_order` which is DPC++ only + sycl::queue myQueue{sycl::property::queue::in_order{}}; + return 0; +} diff --git a/hip/base/executor.hip.cpp b/hip/base/executor.hip.cpp index f41fb69f46c..3e400986e2b 100644 --- a/hip/base/executor.hip.cpp +++ b/hip/base/executor.hip.cpp @@ -138,11 +138,18 @@ void HipExecutor::raw_copy_to(const CudaExecutor *dest, size_type num_bytes, num_bytes)); } #else - GKO_NOT_SUPPORTED(this); + GKO_NOT_SUPPORTED(dest); #endif } +void HipExecutor::raw_copy_to(const DpcppExecutor *dest, size_type num_bytes, + const void *src_ptr, void *dest_ptr) const +{ + GKO_NOT_SUPPORTED(dest); +} + + void HipExecutor::raw_copy_to(const HipExecutor *dest, size_type num_bytes, const void *src_ptr, void *dest_ptr) const { diff --git a/hip/test/base/hip_executor.hip.cpp b/hip/test/base/hip_executor.hip.cpp index 635639fc21e..e0003e91c4b 100644 --- a/hip/test/base/hip_executor.hip.cpp +++ b/hip/test/base/hip_executor.hip.cpp @@ -75,6 +75,11 @@ class ExampleOperation : public gko::Operation { value = -3; } + void run(std::shared_ptr) const override + { + value = -4; + } + void run(std::shared_ptr) const override { hipGetDevice(&value); @@ -248,7 +253,7 @@ TEST_F(HipExecutor, CopiesDataFromHipToHip) omp->copy_from(hip2.get(), 2, copy_hip2, copy); EXPECT_EQ(3, copy[0]); ASSERT_EQ(8, copy[1]); - hip->free(copy_hip2); + hip2->free(copy_hip2); hip->free(orig); } diff --git a/include/ginkgo/core/base/executor.hpp b/include/ginkgo/core/base/executor.hpp index 1df29abc59c..770ad201f15 100644 --- a/include/ginkgo/core/base/executor.hpp +++ b/include/ginkgo/core/base/executor.hpp @@ -37,6 +37,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include #include @@ -46,6 +47,15 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +inline namespace cl { +namespace sycl { + +class queue; + +} // namespace sycl +} // namespace cl + + struct cublasContext; struct cusparseContext; @@ -114,6 +124,9 @@ class ExecutorBase; * void run(const gko::HipExecutor *exec) const override * { os_ << "HIP(" << exec->get_device_id() << ")"; } * + * void run(const gko::DpcppExecutor *exec) const override + * { os_ << "DPC++(" << exec->get_device_id() << ")"; } + * * // This is optional, if not overloaded, defaults to OmpExecutor overload * void run(const gko::ReferenceExecutor *) const override * { os_ << "Reference CPU"; } @@ -142,6 +155,7 @@ class ExecutorBase; * std::cout << *omp << std::endl * << *gko::CudaExecutor::create(0, omp) << std::endl * << *gko::HipExecutor::create(0, omp) << std::endl + * << *gko::DpcppExecutor::create(0, omp) << std::endl * << *gko::ReferenceExecutor::create() << std::endl; * ``` * @@ -151,15 +165,16 @@ class ExecutorBase; * OMP * CUDA(0) * HIP(0) + * DPC++(0) * Reference CPU * ``` * * One might feel that this code is too complicated for such a simple task. * Luckily, there is an overload of the Executor::run() method, which is * designed to facilitate writing simple operations like this one. The method - * takes three closures as input: one which is run for OMP, one for - * CUDA executors, and the last one for HIP executors. Using this method, there - * is no need to implement an Operation subclass: + * takes three closures as input: one which is run for OMP, one for CUDA + * executors, one for HIP executors, and the last one for DPC++ executors. Using + * this method, there is no need to implement an Operation subclass: * * ``` * std::ostream& operator<<(std::ostream &os, const gko::Executor &exec) @@ -174,6 +189,10 @@ class ExecutorBase; * << static_cast(exec) * .get_device_id() * << ")"; }); + * [&]() { os << "DPC++(" // DPC++ closure + * << static_cast(exec) + * .get_device_id() + * << ")"; }); * return os; * } * ``` @@ -250,7 +269,7 @@ private: \ * kernel when the operation is executed. * * The kernels used to bind the operation are searched in `kernels::DEV_TYPE` - * namespace, where `DEV_TYPE` is replaced by `omp`, `cuda`, `hip` and + * namespace, where `DEV_TYPE` is replaced by `omp`, `cuda`, `hip`, `dpcpp` and * `reference`. * * @param _name operation name @@ -278,6 +297,11 @@ private: \ * // hip code * } * } + * namespace dpcpp { + * void my_kernel(int x) { + * // dpcpp code + * } + * } * namespace reference { * void my_kernel(int x) { * // reference code @@ -292,6 +316,7 @@ private: \ * auto omp = OmpExecutor::create(); * auto cuda = CudaExecutor::create(omp, 0); * auto hip = HipExecutor::create(omp, 0); + * auto dpcpp = DpcppExecutor::create(omp, 0); * auto ref = ReferenceExecutor::create(); * * // create the operation @@ -300,50 +325,52 @@ private: \ * omp->run(op); // run omp kernel * cuda->run(op); // run cuda kernel * hip->run(op); // run hip kernel + * dpcpp->run(op); // run DPC++ kernel * ref->run(op); // run reference kernel * } * ``` * * @ingroup Executor */ -#define GKO_REGISTER_OPERATION(_name, _kernel) \ - template \ - class _name##_operation : public Operation { \ - using counts = \ - ::gko::syn::as_list<::gko::syn::range<0, sizeof...(Args)>>; \ - \ - public: \ - explicit _name##_operation(Args &&... args) \ - : data(std::forward(args)...) \ - {} \ - \ - const char *get_name() const noexcept override \ - { \ - static auto name = [this] { \ - std::ostringstream oss; \ - oss << #_kernel << '#' << sizeof...(Args); \ - return oss.str(); \ - }(); \ - return name.c_str(); \ - } \ - \ - GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(OmpExecutor, omp, _kernel); \ - GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(CudaExecutor, cuda, _kernel); \ - GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(HipExecutor, hip, _kernel); \ - GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(ReferenceExecutor, reference, \ - _kernel); \ - \ - private: \ - mutable std::tuple data; \ - }; \ - \ - template \ - static _name##_operation make_##_name(Args &&... args) \ - { \ - return _name##_operation(std::forward(args)...); \ - } \ - static_assert(true, \ - "This assert is used to counter the false positive extra " \ +#define GKO_REGISTER_OPERATION(_name, _kernel) \ + template \ + class _name##_operation : public Operation { \ + using counts = \ + ::gko::syn::as_list<::gko::syn::range<0, sizeof...(Args)>>; \ + \ + public: \ + explicit _name##_operation(Args &&... args) \ + : data(std::forward(args)...) \ + {} \ + \ + const char *get_name() const noexcept override \ + { \ + static auto name = [this] { \ + std::ostringstream oss; \ + oss << #_kernel << '#' << sizeof...(Args); \ + return oss.str(); \ + }(); \ + return name.c_str(); \ + } \ + \ + GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(OmpExecutor, omp, _kernel); \ + GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(CudaExecutor, cuda, _kernel); \ + GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(HipExecutor, hip, _kernel); \ + GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(DpcppExecutor, dpcpp, _kernel); \ + GKO_KERNEL_DETAIL_DEFINE_RUN_OVERLOAD(ReferenceExecutor, reference, \ + _kernel); \ + \ + private: \ + mutable std::tuple data; \ + }; \ + \ + template \ + static _name##_operation make_##_name(Args &&... args) \ + { \ + return _name##_operation(std::forward(args)...); \ + } \ + static_assert(true, \ + "This assert is used to counter the false positive extra " \ "semi-colon warnings") @@ -359,6 +386,8 @@ private: \ * operations executed on the NVIDIA GPU accelerator; * + HipExecutor specifies that the data should be stored and the * operations executed on either an NVIDIA or AMD GPU accelerator; + * + DpcppExecutor specifies that the data should be stored and the + * operations executed on an hardware supporting DPC++; * + ReferenceExecutor executes a non-optimized reference implementation, * which can be used to debug the library. * @@ -436,6 +465,12 @@ class Executor : public log::EnableLogging { template friend class detail::ExecutorBase; +#define GKO_DECLARE_EXECUTOR_FRIEND(_type, ...) friend class _type + + GKO_ENABLE_FOR_ALL_EXECUTORS(GKO_DECLARE_EXECUTOR_FRIEND); + +#undef GKO_DECLARE_EXECUTOR_FRIEND + public: virtual ~Executor() = default; @@ -464,12 +499,13 @@ class Executor : public log::EnableLogging { * @param op_cuda functor to run in case of a CudaExecutor * @param op_hip functor to run in case of a HipExecutor */ - template + template void run(const ClosureOmp &op_omp, const ClosureCuda &op_cuda, - const ClosureHip &op_hip) const + const ClosureHip &op_hip, const ClosureDpcpp &op_dpcpp) const { - LambdaOperation op(op_omp, op_cuda, - op_hip); + LambdaOperation op( + op_omp, op_cuda, op_hip, op_dpcpp); this->run(op); } @@ -530,7 +566,20 @@ class Executor : public log::EnableLogging { this->template log( src_exec, this, reinterpret_cast(src_ptr), reinterpret_cast(dest_ptr), num_elems * sizeof(T)); - this->raw_copy_from(src_exec, num_elems * sizeof(T), src_ptr, dest_ptr); + try { + this->raw_copy_from(src_exec, num_elems * sizeof(T), src_ptr, + dest_ptr); + } catch (NotSupported &err) { + // Unoptimized copy. Try to go through the masters. + auto src_master = src_exec->get_master().get(); + if (num_elems > 0 && src_master != src_exec) { + auto *master_ptr = src_exec->get_master()->alloc(num_elems); + src_master->copy_from(src_exec, num_elems, src_ptr, + master_ptr); + this->copy_from(src_master, num_elems, master_ptr, dest_ptr); + src_master->free(master_ptr); + } + } this->template log( src_exec, this, reinterpret_cast(src_ptr), reinterpret_cast(dest_ptr), num_elems * sizeof(T)); @@ -650,8 +699,10 @@ class Executor : public log::EnableLogging { * @tparam ClosureOmp the type of the first functor * @tparam ClosureCuda the type of the second functor * @tparam ClosureHip the type of the third functor + * @tparam ClosureDpcpp the type of the fourth functor */ - template + template class LambdaOperation : public Operation { public: /** @@ -661,10 +712,15 @@ class Executor : public log::EnableLogging { * and ReferenceExecutor * @param op_cuda a functor object which will be called by CudaExecutor * @param op_hip a functor object which will be called by HipExecutor + * @param op_dpcpp a functor object which will be called by + * DpcppExecutor */ LambdaOperation(const ClosureOmp &op_omp, const ClosureCuda &op_cuda, - const ClosureHip &op_hip) - : op_omp_(op_omp), op_cuda_(op_cuda), op_hip_(op_hip) + const ClosureHip &op_hip, const ClosureDpcpp &op_dpcpp) + : op_omp_(op_omp), + op_cuda_(op_cuda), + op_hip_(op_hip), + op_dpcpp_(op_dpcpp) {} void run(std::shared_ptr) const override @@ -682,10 +738,16 @@ class Executor : public log::EnableLogging { op_hip_(); } + void run(std::shared_ptr) const override + { + op_dpcpp_(); + } + private: ClosureOmp op_omp_; ClosureCuda op_cuda_; ClosureHip op_hip_; + ClosureDpcpp op_dpcpp_; }; }; @@ -755,6 +817,12 @@ namespace detail { template class ExecutorBase : public Executor { +#define GKO_DECLARE_EXECUTOR_FRIEND(_type, ...) friend class _type + + GKO_ENABLE_FOR_ALL_EXECUTORS(GKO_DECLARE_EXECUTOR_FRIEND); + +#undef GKO_DECLARE_EXECUTOR_FRIEND + public: void run(const Operation &op) const override { @@ -1238,6 +1306,122 @@ using DefaultExecutor = HipExecutor; } // namespace kernels +/** + * This is the Executor subclass which represents a DPC++ enhanced device. + * + * @ingroup exec_dpcpp + * @ingroup Executor + */ +class DpcppExecutor : public detail::ExecutorBase, + public std::enable_shared_from_this { + friend class detail::ExecutorBase; + +public: + /** + * Creates a new DpcppExecutor. + * + * @param device_id the DPCPP device id of this device + * @param master an executor on the host that is used to invoke the device + * kernels + * @param device_type a string representing the type of device to consider + * (accelerator, cpu, gpu or all). + */ + static std::shared_ptr create( + int device_id, std::shared_ptr master, + std::string device_type = "gpu"); + + std::shared_ptr get_master() noexcept override; + + std::shared_ptr get_master() const noexcept override; + + void synchronize() const override; + + void run(const Operation &op) const override; + + /** + * Get the DPCPP device id of the device associated to this executor. + */ + int get_device_id() const noexcept { return device_id_; } + + ::cl::sycl::queue *get_queue() const { return queue_.get(); } + + /** + * Get the number of devices present on the system. + */ + static int get_num_devices(std::string device_type); + + /** + * Get the available subgroup sizes for this device. + */ + const std::vector &get_subgroup_sizes() const noexcept + { + return subgroup_sizes_; + } + + /** + * Get the number of Computing Units of this executor. + */ + size_type get_num_computing_units() const noexcept + { + return num_computing_units_; + } + + /** + * Get the maximum work item sizes. + */ + const std::vector &get_max_workitem_sizes() const noexcept + { + return max_workitem_sizes_; + } + + /** + * Get the maximum workgroup size. + */ + size_type get_max_workgroup_size() const noexcept + { + return max_workgroup_size_; + } + +protected: + void set_gpu_property(); + + DpcppExecutor(int device_id, std::shared_ptr master, + std::string device_type = "gpu") + : device_id_(device_id), master_(master), device_type_(device_type) + { + assert(device_id < DpcppExecutor::get_num_devices(device_type_)); + + this->set_gpu_property(); + } + + void *raw_alloc(size_type size) const override; + + void raw_free(void *ptr) const noexcept override; + + GKO_ENABLE_FOR_ALL_EXECUTORS(GKO_OVERRIDE_RAW_COPY_TO); + +private: + int device_id_; + std::shared_ptr master_; + std::string device_type_; + int num_computing_units_{}; + std::vector subgroup_sizes_{}; + std::vector max_workitem_sizes_{}; + size_type max_workgroup_size_{}; + + template + using queue_manager = std::unique_ptr>; + queue_manager<::cl::sycl::queue> queue_; +}; + + +namespace kernels { +namespace dpcpp { +using DefaultExecutor = DpcppExecutor; +} // namespace dpcpp +} // namespace kernels + + #undef GKO_OVERRIDE_RAW_COPY_TO diff --git a/include/ginkgo/core/base/math.hpp b/include/ginkgo/core/base/math.hpp index 0bb6ac9b5a9..9217e00f800 100644 --- a/include/ginkgo/core/base/math.hpp +++ b/include/ginkgo/core/base/math.hpp @@ -41,6 +41,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#ifdef CL_SYCL_LANGUAGE_VERSION +#include +#endif + + #include #include #include @@ -908,7 +913,11 @@ GKO_INLINE GKO_ATTRIBUTES constexpr xstd::enable_if_t::value, remove_complex> abs(const T &x) { +#ifdef CL_SYCL_LANGUAGE_VERSION + return cl::sycl::sqrt(real(x) * real(x) + imag(x) * imag(x)); +#else return sqrt(squared_norm(x)); +#endif } @@ -966,7 +975,11 @@ GKO_INLINE GKO_ATTRIBUTES std::enable_if_t::value, bool> is_finite(const T &value) { constexpr T infinity{detail::infinity_impl::value}; +#ifdef CL_SYCL_LANGUAGE_VERSION + return ::gko::abs(value) < infinity; +#else return abs(value) < infinity; +#endif } @@ -989,6 +1002,25 @@ is_finite(const T &value) } +namespace kernels { +namespace dpcpp { + + +// For now this seems to be useless. Somehow, DPC++ doesn't use this +// declaration and anyway always replace calls to `abs` by `std::abs`. To +// reference this declaration, use `dpcpp::abs`. +using ::gko::abs; + + +#ifdef CL_SYCL_LANGUAGE_VERSION +using cl::sycl::sqrt; +#endif + + +} // namespace dpcpp +} // namespace kernels + + } // namespace gko diff --git a/include/ginkgo/core/base/types.hpp b/include/ginkgo/core/base/types.hpp index 10ef4d5d4cf..d63623fea40 100644 --- a/include/ginkgo/core/base/types.hpp +++ b/include/ginkgo/core/base/types.hpp @@ -392,6 +392,7 @@ GKO_ATTRIBUTES constexpr bool operator!=(precision_reduction x, #define GKO_ENABLE_FOR_ALL_EXECUTORS(_enable_macro) \ _enable_macro(OmpExecutor, omp); \ _enable_macro(HipExecutor, hip); \ + _enable_macro(DpcppExecutor, dpcpp); \ _enable_macro(CudaExecutor, cuda) diff --git a/include/ginkgo/core/base/version.hpp b/include/ginkgo/core/base/version.hpp index 52731ab56e3..33ed03a8872 100644 --- a/include/ginkgo/core/base/version.hpp +++ b/include/ginkgo/core/base/version.hpp @@ -212,6 +212,13 @@ class version_info { */ version hip_version; + /** + * Contains version information of the DPC++ module. + * + * This is the version of the static/shared library called "ginkgo_dpcpp". + */ + version dpcpp_version; + private: static constexpr version get_header_version() noexcept { @@ -229,13 +236,16 @@ class version_info { static version get_hip_version() noexcept; + static version get_dpcpp_version() noexcept; + version_info() : header_version{get_header_version()}, core_version{get_core_version()}, reference_version{get_reference_version()}, omp_version{get_omp_version()}, cuda_version{get_cuda_version()}, - hip_version{get_hip_version()} + hip_version{get_hip_version()}, + dpcpp_version{get_dpcpp_version()} {} }; diff --git a/omp/CMakeLists.txt b/omp/CMakeLists.txt index bcfde819760..6578bd31d5e 100644 --- a/omp/CMakeLists.txt +++ b/omp/CMakeLists.txt @@ -49,6 +49,8 @@ target_compile_options(ginkgo_omp PRIVATE "${GINKGO_COMPILER_FLAGS}") target_link_libraries(ginkgo_omp PUBLIC ginkgo_cuda) # Need to link against ginkgo_hip for the `raw_copy_to(HipExecutor ...)` method target_link_libraries(ginkgo_omp PUBLIC ginkgo_hip) +# Need to link against ginkgo_dpcpp for the `raw_copy_to(DpcppExecutor ...)` method +target_link_libraries(ginkgo_omp PUBLIC ginkgo_dpcpp) ginkgo_default_includes(ginkgo_omp) ginkgo_install_library(ginkgo_omp omp)