[SYCL] :

  * Modifying TensorDeviceSYCL to use `EIGEN_THROW_X`.
  * Modifying TensorMacro to use `EIGEN_TRY/CATCH(X)` macro.
  * Modifying TensorReverse.h to use `EIGEN_DEVICE_REF` instead of `&`.
  * Fixing the SYCL device macro in SpecialFunctionsImpl.h.

4 files changed, 21 insertions, 33 deletions

diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h
index 93efe2f82..6f8b6f193 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h
@@ -16,6 +16,7 @@
 #define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H
 #include <unordered_set>
 
+
 namespace Eigen {
 
 namespace TensorSycl {
@@ -676,21 +677,12 @@ class QueueInterface {
     }
   }
 
-  bool sycl_async_handler(cl::sycl::exception_list l) const {
+  bool sycl_async_handler(cl::sycl::exception_list exceptions) const {
     bool exception_caught = false;
-    for (const auto &e : l) {
+    for (const auto &e : exceptions) {
       if (e) {
         exception_caught = true;
-#ifdef EIGEN_EXCEPTIONS
-        try {
-          std::rethrow_exception(e);
-        } catch (const cl::sycl::exception &e) {
-          std::cerr << e.what() << std::endl;
-        }
-#else
-        std::cerr << "Error detected inside Sycl device." << std::endl;
-        abort();
-#endif
+        EIGEN_THROW_X(e);
       }
     }
     return exception_caught;
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorMacros.h b/unsupported/Eigen/CXX11/src/Tensor/TensorMacros.h
index 3d859f42d..af9e5db70 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorMacros.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorMacros.h
@@ -67,20 +67,16 @@
 #endif
 
 // Define a macro for catching SYCL exceptions if exceptions are enabled
-#if defined(EIGEN_EXCEPTIONS)
-  #define EIGEN_SYCL_TRY_CATCH(X) \
-    do { \
-      try { X; } \
-      catch(const cl::sycl::exception& e) { \
-        std::cerr << "SYCL exception at " \
-                  << __FILE__ << ":" << __LINE__ << std::endl \
-                  << e.what() << std::endl; \
-        std::rethrow_exception(std::current_exception()); \
-      } \
-    } while (false)
-#else
-  #define EIGEN_SYCL_TRY_CATCH(X) X
-#endif
+#define EIGEN_SYCL_TRY_CATCH(X) \
+  do { \
+    EIGEN_TRY {X;} \
+    EIGEN_CATCH(const cl::sycl::exception& e) { \
+      EIGEN_THROW_X(std::runtime_error("SYCL exception at " + \
+                                       std::string(__FILE__) + ":" + \
+                                       std::to_string(__LINE__) + "\n" + \
+                                       e.what())); \
+    } \
+  } while (false)
 
 // Define a macro if local memory flags are unset or one of them is set
 // Setting both flags is the same as unsetting them
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReverse.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReverse.h
index 42205db31..123675196 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorReverse.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReverse.h
@@ -380,7 +380,7 @@ struct TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>, Device
   array<IndexDivisor, NumDims> m_fastStrides;
   TensorEvaluator<ArgType, Device> m_impl;
   ReverseDimensions m_reverse;
-  const Device& m_device;
+  const Device EIGEN_DEVICE_REF m_device;
 };
 
 // Eval as lvalue
diff --git a/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h b/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h
index 5784cbc86..1464a9751 100644
--- a/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h
+++ b/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h
@@ -193,7 +193,7 @@ struct lgamma_impl<float> {
 #if !defined(EIGEN_GPU_COMPILE_PHASE) && (defined(_BSD_SOURCE) || defined(_SVID_SOURCE)) && !defined(__APPLE__) 
     int dummy;
     return ::lgammaf_r(x, &dummy);
-#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
+#elif defined(SYCL_DEVICE_ONLY)
     return cl::sycl::lgamma(x);
 #else
     return ::lgammaf(x);
@@ -208,7 +208,7 @@ struct lgamma_impl<double> {
 #if !defined(EIGEN_GPU_COMPILE_PHASE) && (defined(_BSD_SOURCE) || defined(_SVID_SOURCE)) && !defined(__APPLE__) 
     int dummy;
     return ::lgamma_r(x, &dummy);
-#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
+#elif defined(SYCL_DEVICE_ONLY)
     return cl::sycl::lgamma(x);
 #else
     return ::lgamma(x);
@@ -428,7 +428,7 @@ template <>
 struct erf_impl<float> {
   EIGEN_DEVICE_FUNC
   static EIGEN_STRONG_INLINE float run(float x) {
-#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
+#if defined(SYCL_DEVICE_ONLY)
     return cl::sycl::erf(x);
 #else
     return ::erff(x);
@@ -440,7 +440,7 @@ template <>
 struct erf_impl<double> {
   EIGEN_DEVICE_FUNC
   static EIGEN_STRONG_INLINE double run(double x) {
-#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
+#if defined(SYCL_DEVICE_ONLY)
     return cl::sycl::erf(x);
 #else
     return ::erf(x);
@@ -473,7 +473,7 @@ template <>
 struct erfc_impl<float> {
   EIGEN_DEVICE_FUNC
   static EIGEN_STRONG_INLINE float run(const float x) {
-#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
+#if defined(SYCL_DEVICE_ONLY)
     return cl::sycl::erfc(x);
 #else
     return ::erfcf(x);
@@ -485,7 +485,7 @@ template <>
 struct erfc_impl<double> {
   EIGEN_DEVICE_FUNC
   static EIGEN_STRONG_INLINE double run(const double x) {
-#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
+#if defined(SYCL_DEVICE_ONLY)
     return cl::sycl::erfc(x);
 #else
     return ::erfc(x);