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

5 files changed, 73 insertions, 44 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})
-
-  add_executable(${targetname} ${include_file})
-  add_dependencies(${targetname} ${testname}_integration_header_sycl)
+  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_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 <SYCL/sycl.hpp>
+#include <CL/sycl.hpp>
 #include <iostream>
 #include <map>
 #include <memory>
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<HostExpr>::Type DevExpr;
       typedef typename Eigen::TensorSycl::internal::ConvertToDeviceExpression<LHSHostExpr>::Type LHSDevExpr;
       typedef typename Eigen::TensorSycl::internal::ConvertToDeviceExpression<RHSHostExpr>::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}")