From 052426b824038bbee1c1c48c3df65dfccd79ae24 Mon Sep 17 00:00:00 2001 From: a-doumoulakis Date: Fri, 5 May 2017 19:26:27 +0100 Subject: [PATCH 01/12] 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}") From a5226ce4f77f68a7f8590174ed3cdebb89cbe127 Mon Sep 17 00:00:00 2001 From: a-doumoulakis Date: Wed, 17 May 2017 17:59:30 +0100 Subject: [PATCH 02/12] Add cmake file FindTriSYCL.cmake --- cmake/FindTriSYCL.cmake | 152 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 152 insertions(+) create mode 100644 cmake/FindTriSYCL.cmake diff --git a/cmake/FindTriSYCL.cmake b/cmake/FindTriSYCL.cmake new file mode 100644 index 000000000..cb2154192 --- /dev/null +++ b/cmake/FindTriSYCL.cmake @@ -0,0 +1,152 @@ +#.rst: +# FindTriSYCL +#--------------- +# +# TODO : insert Copyright and licence + +######################### +# FindTriSYCL.cmake +######################### +# +# Tools for finding and building with TriSYCL. +# +# User must define TRISYCL_INCLUDE_DIR pointing to the triSYCL +# include directory. +# +# Latest version of this file can be found at: +# https://github.com/triSYCL/triSYCL + +# Requite CMake version 3.5 or higher +cmake_minimum_required (VERSION 3.5) + +# Check that a supported host compiler can be found +if(CMAKE_COMPILER_IS_GNUCXX) + # Require at least gcc 5.4 + if (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 5.4) + message(FATAL_ERROR + "host compiler - Not found! (gcc version must be at least 5.4)") + else() + message(STATUS "host compiler - gcc ${CMAKE_CXX_COMPILER_VERSION}") + endif() +elseif ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang") + # Require at least clang 3.9 + if (${CMAKE_CXX_COMPILER_VERSION} VERSION_LESS 3.9) + message(FATAL_ERROR + "host compiler - Not found! (clang version must be at least 3.9)") + else() + message(STATUS "host compiler - clang ${CMAKE_CXX_COMPILER_VERSION}") + endif() +else() + message(WARNING + "host compiler - Not found! (triSYCL supports GCC and Clang)") +endif() + +#triSYCL options +option(TRISYCL_OPENMP "triSYCL multi-threading with OpenMP" ON) +option(TRISYCL_OPENCL "triSYCL OpenCL interoperability mode" OFF) +option(TRISYCL_NO_ASYNC "triSYCL use synchronous kernel execution" OFF) +option(TRISYCL_DEBUG "triSCYL use debug mode" OFF) +option(TRISYCL_DEBUG_STRUCTORS "triSYCL trace of object lifetimes" OFF) +option(TRISYCL_TRACE_KERNEL "triSYCL trace of kernel execution" OFF) + +mark_as_advanced(TRISYCL_OPENMP) +mark_as_advanced(TRISYCL_OPENCL) +mark_as_advanced(TRISYCL_NO_ASYNC) +mark_as_advanced(TRISYCL_DEBUG) +mark_as_advanced(TRISYCL_DEBUG_STRUCTORS) +mark_as_advanced(TRISYCL_TRACE_KERNEL) + +#triSYCL definitions +set(CL_SYCL_LANGUAGE_VERSION 220 CACHE VERSION + "Host language version to be used by trisYCL (default is: 220)") +set(TRISYCL_CL_LANGUAGE_VERSION 220 CACHE VERSION + "Device language version to be used by trisYCL (default is: 220)") +#set(TRISYCL_COMPILE_OPTIONS "-std=c++1z -Wall -Wextra") +set(CMAKE_CXX_STANDARD 14) +set(CXX_STANDARD_REQUIRED ON) + + +# Find OpenCL package +if(TRISYCL_OPENCL) + find_package(OpenCL REQUIRED) + if(UNIX) + set(BOOST_COMPUTE_INCPATH /usr/include/compute CACHE PATH + "Path to Boost.Compute headers (default is: /usr/include/compute)") + endif(UNIX) +endif() + +# Find OpenMP package +if(TRISYCL_OPENMP) + find_package(OpenMP REQUIRED) +endif() + +# Find Boost +find_package(Boost 1.58 REQUIRED COMPONENTS chrono log) + +# If debug or trace we need boost log +if(TRISYCL_DEBUG OR TRISYCL_DEBUG_STRUCTORS OR TRISYCL_TRACE_KERNEL) + set(LOG_NEEDED ON) +else() + set(LOG_NEEDED OFF) +endif() + +find_package(Threads REQUIRED) + +# Find triSYCL directory +if(NOT TRISYCL_INCLUDE_DIR) + message(FATAL_ERROR + "triSYCL include directory - Not found! (please set TRISYCL_INCLUDE_DIR") +else() + message(STATUS "triSYCL include directory - Found ${TRISYCL_INCLUDE_DIR}") +endif() + +####################### +# add_sycl_to_target +####################### +# +# Sets the proper flags and includes for the target compilation. +# +# targetName : Name of the target to add a SYCL to. +# sourceFile : Source file to be compiled for SYCL. +# binaryDir : Intermediate directory to output the integration header. +# +function(add_sycl_to_target targetName sourceFile binaryDir) + + # Add include directories to the "#include <>" paths + target_include_directories (${targetName} PUBLIC + ${TRISYCL_INCLUDE_DIR} + ${Boost_INCLUDE_DIRS} + $<$:${OpenCL_INCLUDE_DIRS}> + $<$:${BOOST_COMPUTE_INCPATH}>) + + + # Link dependencies + target_link_libraries(${targetName} PUBLIC + $<$:${OpenCL_LIBRARIES}> + Threads::Threads + $<$:Boost::log> + Boost::chrono) + + + # Compile definitions + target_compile_definitions(${targetName} PUBLIC + $<$:TRISYCL_NO_ASYNC> + $<$:TRISYCL_OPENCL> + $<$:TRISYCL_DEBUG> + $<$:TRISYCL_DEBUG_STRUCTORS> + $<$:TRISYCL_TRACE_KERNEL> + $<$:BOOST_LOG_DYN_LINK>) + + # C++ and OpenMP requirements + target_compile_options(${targetName} PUBLIC + ${TRISYCL_COMPILE_OPTIONS} + $<$:${OpenMP_CXX_FLAGS}>) + + if(${TRISYCL_OPENMP} AND (NOT WIN32)) + # Does not support generator expressions + set_target_properties(${targetName} + PROPERTIES + LINK_FLAGS ${OpenMP_CXX_FLAGS}) + endif(${TRISYCL_OPENMP} AND (NOT WIN32)) + +endfunction(add_sycl_to_target) From 61d7f3664a06c64e39a0bbccb610592bcaee92cd Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Mon, 22 May 2017 14:58:28 +0100 Subject: [PATCH 03/12] Fixing Cmake Dependency for SYCL --- cmake/EigenTesting.cmake | 18 +++++++++--------- cmake/FindComputeCpp.cmake | 8 +++++--- unsupported/test/CMakeLists.txt | 2 +- 3 files changed, 15 insertions(+), 13 deletions(-) diff --git a/cmake/EigenTesting.cmake b/cmake/EigenTesting.cmake index a92a2978b..1fcdcc7bb 100644 --- a/cmake/EigenTesting.cmake +++ b/cmake/EigenTesting.cmake @@ -120,23 +120,23 @@ macro(ei_add_test_internal_sycl testname testname_with_suffix) set(filename ${testname}.cpp) endif() - set( include_file ${CMAKE_CURRENT_BINARY_DIR}/inc_${filename}) - set( bc_file ${CMAKE_CURRENT_BINARY_DIR}/${filename}) - set( host_file ${CMAKE_CURRENT_SOURCE_DIR}/${filename}) + set( include_file "${CMAKE_CURRENT_BINARY_DIR}/inc_${filename}") + set( bc_file "${CMAKE_CURRENT_BINARY_DIR}/${filename}.sycl") + 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 + COMMAND ${CMAKE_COMMAND} -E echo "\\#include \\\"${bc_file}\\\"" >> ${include_file} + DEPENDS ${host_file} ${bc_file} 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_custom_target("${testname}_integration_header_sycl" DEPENDS ${include_file}) add_executable(${targetname} ${include_file}) - add_dependencies(${targetname} ${testname}_integration_header_sycl) - add_sycl_to_target(${targetname} ${filename} ${CMAKE_CURRENT_BINARY_DIR}) + add_dependencies(${targetname} "${testname}_integration_header_sycl") + add_sycl_to_target(${targetname} ${CMAKE_CURRENT_BINARY_DIR} ${filename}) if (targetname MATCHES "^eigen2_") add_dependencies(eigen2_buildtests ${targetname}) @@ -720,4 +720,4 @@ macro(ei_test_get_compilerver_from_cxx_version_string) ei_test1_get_compilerver_from_cxx_version_string("i686-apple-darwin11-llvm-g++-4.2 (GCC) 4.2.1 (Based on Apple Inc. build 5658) (LLVM build 2335.15.00)" "llvm-g++" "4.2.1") ei_test1_get_compilerver_from_cxx_version_string("g++-mp-4.4 (GCC) 4.4.6" "g++" "4.4.6") ei_test1_get_compilerver_from_cxx_version_string("g++-mp-4.4 (GCC) 2011" "g++" "4.4") -endmacro(ei_test_get_compilerver_from_cxx_version_string) +endmacro(ei_test_get_compilerver_from_cxx_version_string) \ No newline at end of file diff --git a/cmake/FindComputeCpp.cmake b/cmake/FindComputeCpp.cmake index 27e5c9b1f..90222ed0e 100644 --- a/cmake/FindComputeCpp.cmake +++ b/cmake/FindComputeCpp.cmake @@ -201,7 +201,8 @@ function(__build_spir targetName sourceFile binaryDir) ${device_compiler_includes} -o ${outputSyclFile} -c ${CMAKE_CURRENT_SOURCE_DIR}/${sourceFile} - DEPENDS ${sourceFile} + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/${sourceFile} + IMPLICIT_DEPENDS CXX "${CMAKE_CURRENT_SOURCE_DIR}/${sourceFile}" WORKING_DIRECTORY ${binaryDir} COMMENT "Building ComputeCpp integration header file ${outputSyclFile}") @@ -233,8 +234,9 @@ endfunction() # sourceFile : Source file to be compiled for SYCL. # binaryDir : Intermediate directory to output the integration header. # -function(add_sycl_to_target targetName sourceFile binaryDir) +function(add_sycl_to_target targetName binaryDir sourceFile) + set(sourceFiles ${sourceFiles} ${ARGN}) # Add custom target to run compute++ and generate the integration header __build_spir(${targetName} ${sourceFile} ${binaryDir}) @@ -242,4 +244,4 @@ function(add_sycl_to_target targetName sourceFile binaryDir) target_link_libraries(${targetName} PUBLIC ${COMPUTECPP_RUNTIME_LIBRARY} PUBLIC ${OpenCL_LIBRARIES}) -endfunction(add_sycl_to_target) +endfunction(add_sycl_to_target) \ No newline at end of file diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt index cdf151f15..196a432ff 100644 --- a/unsupported/test/CMakeLists.txt +++ b/unsupported/test/CMakeLists.txt @@ -172,7 +172,7 @@ if(EIGEN_TEST_CXX11) 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_volume_patch_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") endif(EIGEN_TEST_SYCL) From 2d17128d6f885ab839f2ecde360268f5b523426c Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Mon, 22 May 2017 16:40:33 +0100 Subject: [PATCH 04/12] Fixing suported device list. --- .../Eigen/CXX11/src/Tensor/TensorDeviceSycl.h | 34 +++++++++---------- 1 file changed, 16 insertions(+), 18 deletions(-) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index c5142b7c9..06df1030b 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -81,28 +81,26 @@ struct memsetCghFunctor{ } }; - //get_devices returns all the available opencl devices. Either use device_selector or exclude devices that computecpp does not support (AMD OpenCL for CPU and intel GPU) +//get_devices returns all the available opencl devices. Either use device_selector or exclude devices that computecpp does not support (AMD OpenCL for CPU and intel GPU) EIGEN_STRONG_INLINE auto get_sycl_supported_devices()->decltype(cl::sycl::device::get_devices()){ - auto devices = cl::sycl::device::get_devices(); - std::vector::iterator it =devices.begin(); - while(it!=devices.end()) { - ///FIXME: Currently there is a bug in amd cpu OpenCL - auto name = (*it).template get_info(); - std::transform(name.begin(), name.end(), name.begin(), ::tolower); - auto vendor = (*it).template get_info(); +std::vector supported_devices; +auto plafrom_list =cl::sycl::platform::get_platforms(); +for(const auto& platform : plafrom_list){ + auto device_list = platform.get_devices(); + auto platform_name =platform.template get_info(); + std::transform(platform_name.begin(), platform_name.end(), platform_name.begin(), ::tolower); + for(const auto& device : device_list){ + auto vendor = device.template get_info(); std::transform(vendor.begin(), vendor.end(), vendor.begin(), ::tolower); - - if((*it).is_cpu() && vendor.find("amd")!=std::string::npos && vendor.find("apu") == std::string::npos){ // remove amd cpu as it is not supported by computecpp allow APUs - it = devices.erase(it); - //FIXME: currently there is a bug in intel gpu driver regarding memory allignment issue. - }else if((*it).is_gpu() && name.find("intel")!=std::string::npos){ - it = devices.erase(it); - } - else{ - ++it; + bool unsuported_condition = (device.is_cpu() && platform_name.find("amd")!=std::string::npos && vendor.find("apu") == std::string::npos) || + (device.is_gpu() && platform_name.find("intel")!=std::string::npos); + if(!unsuported_condition){ + std::cout << "Platform name "<< platform_name << std::endl; + supported_devices.push_back(device); } } - return devices; +} +return supported_devices; } class QueueInterface { From 76c0fc1f955eda3d243db8960cb6fee9a5305112 Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Mon, 22 May 2017 16:49:32 +0100 Subject: [PATCH 05/12] Fixing SYCL alignment issue required by TensorFlow. --- .../Eigen/CXX11/src/Tensor/TensorDeviceSycl.h | 23 +++++++++++++++++-- 1 file changed, 21 insertions(+), 2 deletions(-) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index c5142b7c9..627c0ab19 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -14,7 +14,23 @@ #if defined(EIGEN_USE_SYCL) && !defined(EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H) #define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H +template struct CheckAlignStatically{ + static const bool Val= (((Align&(Align-1))==0) && (Align >= sizeof(void *))); +}; +template +struct Conditional_Allocate{ +EIGEN_ALWAYS_INLINE static void* conditional_allocate(std::size_t elements){ + return aligned_alloc(Align, elements); +} +}; +template +struct Conditional_Allocate{ + +EIGEN_ALWAYS_INLINE static void* conditional_allocate(std::size_t elements){ + return malloc(elements); +} +}; template > struct SyclAllocator { typedef Scalar value_type; @@ -22,7 +38,10 @@ struct SyclAllocator { typedef typename std::allocator_traits::size_type size_type; SyclAllocator( ){}; - Scalar* allocate(std::size_t elements) { return static_cast(aligned_alloc(Align, elements)); } + Scalar* allocate(std::size_t elements) { + return static_cast(Conditional_Allocate::Val, Align>::conditional_allocate(elements)); + // return static_cast(aligned_alloc(Align, elements)); + } void deallocate(Scalar * p, std::size_t size) { EIGEN_UNUSED_VARIABLE(size); free(p); } }; @@ -533,4 +552,4 @@ struct SyclKernelDevice:DefaultDevice{}; } // end namespace Eigen -#endif // EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H +#endif // EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H \ No newline at end of file From b42d775f1350838234f9f05bf64a4f7c12282218 Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Tue, 23 May 2017 10:51:14 +0100 Subject: [PATCH 06/12] Temporarry branch for synch with upstream From 0cb3c7c7dd46d318921b5e140afc73fded87b068 Mon Sep 17 00:00:00 2001 From: Duncan McBain Date: Wed, 24 May 2017 12:24:21 +0100 Subject: [PATCH 07/12] Update FindComputeCpp.cmake with new changes from SDK --- cmake/FindComputeCpp.cmake | 49 ++++++++++++++++++++++++++------------ 1 file changed, 34 insertions(+), 15 deletions(-) diff --git a/cmake/FindComputeCpp.cmake b/cmake/FindComputeCpp.cmake index 90222ed0e..e61dedc46 100644 --- a/cmake/FindComputeCpp.cmake +++ b/cmake/FindComputeCpp.cmake @@ -38,11 +38,6 @@ if(CMAKE_COMPILER_IS_GNUCXX) if (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 4.8) message(FATAL_ERROR "host compiler - Not found! (gcc version must be at least 4.8)") - # Require the GCC dual ABI to be disabled for 5.1 or higher - elseif (CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 5.1) - set(COMPUTECPP_DISABLE_GCC_DUAL_ABI "True") - message(STATUS - "host compiler - gcc ${CMAKE_CXX_COMPILER_VERSION} (note pre 5.1 gcc ABI enabled)") else() message(STATUS "host compiler - gcc ${CMAKE_CXX_COMPILER_VERSION}") endif() @@ -64,6 +59,12 @@ option(COMPUTECPP_64_BIT_CODE "Compile device code in 64 bit mode" ${COMPUTECPP_64_BIT_DEFAULT}) mark_as_advanced(COMPUTECPP_64_BIT_CODE) +option(COMPUTECPP_DISABLE_GCC_DUAL_ABI "Compile with pre-5.1 ABI" OFF) +mark_as_advanced(COMPUTECPP_DISABLE_GCC_DUAL_ABI) + +set(COMPUTECPP_USER_FLAGS "" CACHE STRING "User flags for compute++") +mark_as_advanced(COMPUTECPP_USER_FLAGS) + # Find OpenCL package find_package(OpenCL REQUIRED) @@ -74,7 +75,6 @@ if(NOT COMPUTECPP_PACKAGE_ROOT_DIR) else() message(STATUS "ComputeCpp package - Found") endif() -option(COMPUTECPP_PACKAGE_ROOT_DIR "Path to the ComputeCpp Package") # Obtain the path to compute++ find_program(COMPUTECPP_DEVICE_COMPILER compute++ PATHS @@ -138,8 +138,6 @@ else() message(STATUS "compute++ flags - ${COMPUTECPP_DEVICE_COMPILER_FLAGS}") endif() -set(COMPUTECPP_DEVICE_COMPILER_FLAGS ${COMPUTECPP_DEVICE_COMPILER_FLAGS} -sycl-compress-name -Wall -no-serial-memop -DEIGEN_NO_ASSERTION_CHECKING=1) - # Check if the platform is supported execute_process(COMMAND ${COMPUTECPP_INFO_TOOL} "--dump-is-supported" OUTPUT_VARIABLE COMPUTECPP_PLATFORM_IS_SUPPORTED @@ -155,6 +153,13 @@ else() endif() endif() +set(COMPUTECPP_USER_FLAGS + -sycl-compress-name + -Wall + -no-serial-memop + -DEIGEN_NO_ASSERTION_CHECKING=1 + ) + #################### # __build_sycl #################### @@ -165,8 +170,11 @@ endif() # targetName : Name of the target. # sourceFile : Source file to be compiled. # binaryDir : Intermediate directory to output the integration header. +# fileCounter : Counter included in name of custom target. Different counter +# values prevent duplicated names of custom target when source files with the same name, +# but located in different directories, are used for the same target. # -function(__build_spir targetName sourceFile binaryDir) +function(__build_spir targetName sourceFile binaryDir fileCounter) # Retrieve source file name. get_filename_component(sourceFileName ${sourceFile} NAME) @@ -175,12 +183,16 @@ function(__build_spir targetName sourceFile binaryDir) set(outputSyclFile ${binaryDir}/${sourceFileName}.sycl) # Add any user-defined include to the device compiler + set(device_compiler_includes "") get_property(includeDirectories DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} PROPERTY INCLUDE_DIRECTORIES) - set(device_compiler_includes "") foreach(directory ${includeDirectories}) set(device_compiler_includes "-I${directory}" ${device_compiler_includes}) endforeach() + get_target_property(targetIncludeDirectories ${targetName} INCLUDE_DIRECTORIES) + foreach(directory ${targetIncludeDirectories}) + set(device_compiler_includes "-I${directory}" ${device_compiler_includes}) + endforeach() if (CMAKE_INCLUDE_PATH) foreach(directory ${CMAKE_INCLUDE_PATH}) set(device_compiler_includes "-I${directory}" @@ -188,6 +200,9 @@ function(__build_spir targetName sourceFile binaryDir) endforeach() endif() + set(COMPUTECPP_DEVICE_COMPILER_FLAGS + ${COMPUTECPP_DEVICE_COMPILER_FLAGS} + ${COMPUTECPP_USER_FLAGS}) # Convert argument list format separate_arguments(COMPUTECPP_DEVICE_COMPILER_FLAGS) @@ -204,7 +219,7 @@ function(__build_spir targetName sourceFile binaryDir) DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/${sourceFile} IMPLICIT_DEPENDS CXX "${CMAKE_CURRENT_SOURCE_DIR}/${sourceFile}" WORKING_DIRECTORY ${binaryDir} - COMMENT "Building ComputeCpp integration header file ${outputSyclFile}") + COMMENT "Building ComputeCpp integration header file ${outputSyclFile}") # Add a custom target for the generated integration header add_custom_target(${targetName}_integration_header DEPENDS ${outputSyclFile}) @@ -231,17 +246,21 @@ endfunction() # target and sets a dependancy on that new command. # # targetName : Name of the target to add a SYCL to. -# sourceFile : Source file to be compiled for SYCL. # binaryDir : Intermediate directory to output the integration header. +# sourceFiles : Source files to be compiled for SYCL. # -function(add_sycl_to_target targetName binaryDir sourceFile) +function(add_sycl_to_target targetName binaryDir sourceFiles) set(sourceFiles ${sourceFiles} ${ARGN}) + set(fileCounter 0) # Add custom target to run compute++ and generate the integration header - __build_spir(${targetName} ${sourceFile} ${binaryDir}) + foreach(sourceFile ${sourceFiles}) + __build_spir(${targetName} ${sourceFile} ${binaryDir} ${fileCounter}) + MATH(EXPR fileCounter "${fileCounter} + 1") + endforeach() # Link with the ComputeCpp runtime library target_link_libraries(${targetName} PUBLIC ${COMPUTECPP_RUNTIME_LIBRARY} PUBLIC ${OpenCL_LIBRARIES}) -endfunction(add_sycl_to_target) \ No newline at end of file +endfunction(add_sycl_to_target) From 9ef5c948ba0ae4e24941b1319aac0b7584f10795 Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Wed, 24 May 2017 13:11:16 +0100 Subject: [PATCH 08/12] Fixing Cmake for gcc>=5. From fb853a857a5d3e06687c6736084c20e984b7347d Mon Sep 17 00:00:00 2001 From: a-doumoulakis Date: Wed, 24 May 2017 17:50:15 +0100 Subject: [PATCH 09/12] Restore misplaced comment --- unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index 3105cc122..06df1030b 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -81,7 +81,7 @@ struct memsetCghFunctor{ } }; - //get_devices returns all the available opencl devices. Either use device_selector or exclude devices that computecpp does not support (AMD OpenCL for CPU and intel GPU) +//get_devices returns all the available opencl devices. Either use device_selector or exclude devices that computecpp does not support (AMD OpenCL for CPU and intel GPU) EIGEN_STRONG_INLINE auto get_sycl_supported_devices()->decltype(cl::sycl::device::get_devices()){ std::vector supported_devices; auto plafrom_list =cl::sycl::platform::get_platforms(); From e3f964ed55a96d0c94814d07ab88d8805e0c2eec Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Thu, 25 May 2017 11:17:26 +0100 Subject: [PATCH 10/12] Applying Benoit's comment;removing dead code. --- unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h | 1 - 1 file changed, 1 deletion(-) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index 627c0ab19..27c10a0cf 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -40,7 +40,6 @@ struct SyclAllocator { SyclAllocator( ){}; Scalar* allocate(std::size_t elements) { return static_cast(Conditional_Allocate::Val, Align>::conditional_allocate(elements)); - // return static_cast(aligned_alloc(Align, elements)); } void deallocate(Scalar * p, std::size_t size) { EIGEN_UNUSED_VARIABLE(size); free(p); } }; From c3bd860de8f1d473a05072917f56b380ba59533b Mon Sep 17 00:00:00 2001 From: a-doumoulakis Date: Thu, 25 May 2017 18:46:18 +0100 Subject: [PATCH 11/12] Modification upon request - Remove warning suppression --- CMakeLists.txt | 4 ---- 1 file changed, 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 54f3a7509..2e1648fd1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -151,10 +151,6 @@ 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 From 0370d3576e1df4f898b9030fbc269ec0488d3969 Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Fri, 26 May 2017 16:01:48 +0100 Subject: [PATCH 12/12] Applying Ronnan's comments. --- .../Eigen/CXX11/src/Tensor/TensorDeviceSycl.h | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index 27c10a0cf..09c96317e 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -14,22 +14,22 @@ #if defined(EIGEN_USE_SYCL) && !defined(EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H) #define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H -template struct CheckAlignStatically{ +template struct CheckAlignStatically { static const bool Val= (((Align&(Align-1))==0) && (Align >= sizeof(void *))); }; template -struct Conditional_Allocate{ +struct Conditional_Allocate { -EIGEN_ALWAYS_INLINE static void* conditional_allocate(std::size_t elements){ - return aligned_alloc(Align, elements); -} + EIGEN_ALWAYS_INLINE static void* conditional_allocate(std::size_t elements) { + return aligned_alloc(Align, elements); + } }; template -struct Conditional_Allocate{ +struct Conditional_Allocate { -EIGEN_ALWAYS_INLINE static void* conditional_allocate(std::size_t elements){ - return malloc(elements); -} + EIGEN_ALWAYS_INLINE static void* conditional_allocate(std::size_t elements){ + return malloc(elements); + } }; template > struct SyclAllocator { @@ -551,4 +551,4 @@ struct SyclKernelDevice:DefaultDevice{}; } // end namespace Eigen -#endif // EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H \ No newline at end of file +#endif // EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H