From 052426b824038bbee1c1c48c3df65dfccd79ae24 Mon Sep 17 00:00:00 2001 From: a-doumoulakis Date: Fri, 5 May 2017 19:26:27 +0100 Subject: [PATCH] Add support for triSYCL Eigen is now able to use triSYCL with EIGEN_SYCL_TRISYCL and TRISYCL_INCLUDE_DIR options Fix contraction kernel with correct nd_item dimension --- CMakeLists.txt | 15 ++++- cmake/EigenTesting.cmake | 37 ++++++++---- unsupported/Eigen/CXX11/Tensor | 2 +- .../CXX11/src/Tensor/TensorContractionSycl.h | 2 +- unsupported/test/CMakeLists.txt | 59 +++++++++++-------- 5 files changed, 72 insertions(+), 43 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index fe4227cbb..54f3a7509 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -151,6 +151,10 @@ if(NOT MSVC) ei_add_cxx_compiler_flag("-Wenum-conversion") ei_add_cxx_compiler_flag("-Wc++11-extensions") ei_add_cxx_compiler_flag("-Wdouble-promotion") + ei_add_cxx_compiler_flag("-Wno-unused-parameter") + ei_add_cxx_compiler_flag("-Wno-ignored-attributes") + ei_add_cxx_compiler_flag("-Wno-ignored-qualifiers") + ei_add_cxx_compiler_flag("-Wno-sign-compare") # ei_add_cxx_compiler_flag("-Wconversion") # -Wshadow is insanely too strict with gcc, hopefully it will become usable with gcc 6 @@ -437,10 +441,17 @@ endif() # add SYCL option(EIGEN_TEST_SYCL "Add Sycl support." OFF) +option(EIGEN_SYCL_TRISYCL "Use the triSYCL Sycl implementation (ComputeCPP by default)." OFF) if(EIGEN_TEST_SYCL) set (CMAKE_MODULE_PATH "${CMAKE_ROOT}/Modules" "cmake/Modules/" "${CMAKE_MODULE_PATH}") - include(FindComputeCpp) -endif() + if(EIGEN_SYCL_TRISYCL) + message(STATUS "Using triSYCL") + include(FindTriSYCL) + else(EIGEN_SYCL_TRISYCL) + message(STATUS "Using ComputeCPP SYCL") + include(FindComputeCpp) + endif(EIGEN_SYCL_TRISYCL) +endif(EIGEN_TEST_SYCL) add_subdirectory(unsupported) diff --git a/cmake/EigenTesting.cmake b/cmake/EigenTesting.cmake index a92a2978b..848233342 100644 --- a/cmake/EigenTesting.cmake +++ b/cmake/EigenTesting.cmake @@ -111,7 +111,7 @@ endmacro(ei_add_test_internal) # SYCL macro(ei_add_test_internal_sycl testname testname_with_suffix) - include_directories( SYSTEM ${COMPUTECPP_PACKAGE_ROOT_DIR}/include) + set(targetname ${testname_with_suffix}) if(EIGEN_ADD_TEST_FILENAME_EXTENSION) @@ -124,18 +124,25 @@ macro(ei_add_test_internal_sycl testname testname_with_suffix) set( bc_file ${CMAKE_CURRENT_BINARY_DIR}/${filename}) set( host_file ${CMAKE_CURRENT_SOURCE_DIR}/${filename}) - ADD_CUSTOM_COMMAND( - OUTPUT ${include_file} - COMMAND ${CMAKE_COMMAND} -E echo "\\#include \\\"${host_file}\\\"" > ${include_file} - COMMAND ${CMAKE_COMMAND} -E echo "\\#include \\\"${bc_file}.sycl\\\"" >> ${include_file} - DEPENDS ${filename} ${bc_file}.sycl - COMMENT "Building ComputeCpp integration header file ${include_file}" - ) - # Add a custom target for the generated integration header - add_custom_target(${testname}_integration_header_sycl DEPENDS ${include_file}) + if(NOT EIGEN_SYCL_TRISYCL) + include_directories( SYSTEM ${COMPUTECPP_PACKAGE_ROOT_DIR}/include) + + ADD_CUSTOM_COMMAND( + OUTPUT ${include_file} + COMMAND ${CMAKE_COMMAND} -E echo "\\#include \\\"${host_file}\\\"" > ${include_file} + COMMAND ${CMAKE_COMMAND} -E echo "\\#include \\\"${bc_file}.sycl\\\"" >> ${include_file} + DEPENDS ${filename} ${bc_file}.sycl + COMMENT "Building ComputeCpp integration header file ${include_file}" + ) + + # Add a custom target for the generated integration header + add_custom_target(${testname}_integration_header_sycl DEPENDS ${include_file}) + add_executable(${targetname} ${include_file}) + add_dependencies(${targetname} ${testname}_integration_header_sycl) + else() + add_executable(${targetname} ${host_file}) + endif() - add_executable(${targetname} ${include_file}) - add_dependencies(${targetname} ${testname}_integration_header_sycl) add_sycl_to_target(${targetname} ${filename} ${CMAKE_CURRENT_BINARY_DIR}) if (targetname MATCHES "^eigen2_") @@ -467,7 +474,11 @@ macro(ei_testing_print_summary) endif() if(EIGEN_TEST_SYCL) - message(STATUS "SYCL: ON") + if(EIGEN_SYCL_TRISYCL) + message(STATUS "SYCL: ON (using triSYCL)") + else() + message(STATUS "SYCL: ON (using computeCPP)") + endif() else() message(STATUS "SYCL: OFF") endif() diff --git a/unsupported/Eigen/CXX11/Tensor b/unsupported/Eigen/CXX11/Tensor index 39916092b..5d71a9c25 100644 --- a/unsupported/Eigen/CXX11/Tensor +++ b/unsupported/Eigen/CXX11/Tensor @@ -19,7 +19,7 @@ #undef isnan #undef isinf #undef isfinite -#include +#include #include #include #include diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h index 5b4c3c5bd..e6840bc87 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h @@ -195,7 +195,7 @@ typename HostExpr::Index LocalThreadSizeM, typename HostExpr::Index LocalThreadS m_j_strides(m_j_strides_), m_right_nocontract_strides(m_right_nocontract_strides_), left_tuple_of_accessors(left_tuple_of_accessors_), right_tuple_of_accessors(right_tuple_of_accessors_), dev(dev_){} - void operator()(cl::sycl::nd_item<1> itemID) { + void operator()(cl::sycl::nd_item<2> itemID) { typedef typename Eigen::TensorSycl::internal::ConvertToDeviceExpression::Type DevExpr; typedef typename Eigen::TensorSycl::internal::ConvertToDeviceExpression::Type LHSDevExpr; typedef typename Eigen::TensorSycl::internal::ConvertToDeviceExpression::Type RHSDevExpr; diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt index cdf151f15..9a9124640 100644 --- a/unsupported/test/CMakeLists.txt +++ b/unsupported/test/CMakeLists.txt @@ -152,33 +152,40 @@ endif() if(EIGEN_TEST_CXX11) if(EIGEN_TEST_SYCL) - ei_add_test_sycl(cxx11_tensor_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_forced_eval_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_broadcast_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_device_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_reduction_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_morphing_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_shuffling_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_padding_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_builtins_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_contract_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_concatenation_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_reverse_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_convolution_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_striding_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_chipping_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_layout_swap_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_inflation_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_generator_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_patch_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_image_patch_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_volume_patcP_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_argmax_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_custom_op_sycl "-std=c++11") + if(EIGEN_SYCL_TRISYCL) + set(CMAKE_CXX_STANDARD 14) + set(STD_CXX_FLAG "-std=c++1z") + else(EIGEN_SYCL_TRISYCL) + # It should be safe to always run these tests as there is some fallback code for + # older compiler that don't support cxx11. + set(CMAKE_CXX_STANDARD 11) + set(STD_CXX_FLAG "-std=c++11") + endif(EIGEN_SYCL_TRISYCL) + + ei_add_test_sycl(cxx11_tensor_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_forced_eval_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_broadcast_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_device_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_reduction_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_morphing_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_shuffling_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_padding_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_builtins_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_contract_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_concatenation_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_reverse_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_convolution_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_striding_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_chipping_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_layout_swap_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_inflation_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_generator_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_patch_sycl ${STD_CXX_FLAG}) + #ei_add_test_sycl(cxx11_tensor_image_patch_sycl ${STD_CXX_FLAG}) + #ei_add_test_sycl(cxx11_tensor_volume_patch_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_argmax_sycl ${STD_CXX_FLAG}) + ei_add_test_sycl(cxx11_tensor_custom_op_sycl ${STD_CXX_FLAG}) endif(EIGEN_TEST_SYCL) - # It should be safe to always run these tests as there is some fallback code for - # older compiler that don't support cxx11. - set(CMAKE_CXX_STANDARD 11) ei_add_test(cxx11_eventcount "-pthread" "${CMAKE_THREAD_LIBS_INIT}") ei_add_test(cxx11_runqueue "-pthread" "${CMAKE_THREAD_LIBS_INIT}")