[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch.

* Unifying all loadLocalTile from lhs and rhs to an extract_block function.
* Adding get_tensor operation which was missing in TensorContractionMapper.
* Adding the -D method missing from cmake for Disable_Skinny Contraction operation.
* Wrapping all the indices in TensorScanSycl into Scan parameter struct.
* Fixing typo in Device SYCL
* Unifying load to private register for tall/skinny no shared
* Unifying load to vector tile for tensor-vector/vector-tensor operation
* Removing all the LHS/RHS class for extracting data from global
* Removing Outputfunction from TensorContractionSkinnyNoshared.
* Combining the local memory version of tall/skinny and normal tensor contraction into one kernel.
* Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel.
* Combining General Tensor-Vector and VectorTensor contraction into one kernel.
* Making double buffering optional for Tensor contraction when local memory is version is used.
* Modifying benchmark to accept custom Reduction Sizes
* Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host
* Adding Test for SYCL
* Modifying SYCL CMake

1 files changed, 100 insertions, 33 deletions

diff --git a/bench/tensors/tensor_benchmarks_sycl.cc b/bench/tensors/tensor_benchmarks_sycl.cc
index cb6daac15..b8a096684 100644
--- a/bench/tensors/tensor_benchmarks_sycl.cc
+++ b/bench/tensors/tensor_benchmarks_sycl.cc
@@ -5,19 +5,76 @@
 
 #include "tensor_benchmarks.h"
 
-#define BM_FuncGPU(FUNC)                                                       \
-  static void BM_##FUNC(int iters, int N) {                                    \
-    StopBenchmarkTiming();                                                     \
-    cl::sycl::gpu_selector selector;                                           \
-    Eigen::QueueInterface queue(selector);                                     \
-    Eigen::SyclDevice device(&queue);                                          \
-    BenchmarkSuite<Eigen::SyclDevice, float> suite(device, N);                 \
-    suite.FUNC(iters);                                                         \
-  }                                                                            \
+cl::sycl::gpu_selector selector;
+Eigen::QueueInterface queue(selector);
+#define BM_FuncWithInput2DimsGPU(FUNC, D1, D2)                      \
+  static void BM_##FUNC##_##D1##x##D2(int iters, int N) {           \
+    StopBenchmarkTiming();                                          \
+    Eigen::SyclDevice device(&queue);                               \
+    BenchmarkSuite<Eigen::SyclDevice, float> suite(device, D1, D2); \
+    suite.FUNC(iters);                                              \
+  }                                                                 \
+  BENCHMARK_RANGE(BM_##FUNC##_##D1##x##D2, 10, 10);
+
+BM_FuncWithInput2DimsGPU(rowReduction, 256, 100352);
+BM_FuncWithInput2DimsGPU(rowReduction, 64, 100352);
+BM_FuncWithInput2DimsGPU(rowReduction, 512, 25088);
+BM_FuncWithInput2DimsGPU(rowReduction, 128, 25088);
+BM_FuncWithInput2DimsGPU(rowReduction, 102, 6272);
+BM_FuncWithInput2DimsGPU(rowReduction, 256, 6272);
+BM_FuncWithInput2DimsGPU(rowReduction, 204, 1568);
+BM_FuncWithInput2DimsGPU(rowReduction, 512, 1568);
+BM_FuncWithInput2DimsGPU(rowReduction, 1024, 1568);
+BM_FuncWithInput2DimsGPU(rowReduction, 2048, 1568);
+
+BM_FuncWithInput2DimsGPU(colReduction, 100352, 256);
+BM_FuncWithInput2DimsGPU(colReduction, 100352, 64);
+BM_FuncWithInput2DimsGPU(colReduction, 25088, 512);
+BM_FuncWithInput2DimsGPU(colReduction, 6272, 102);
+BM_FuncWithInput2DimsGPU(colReduction, 25088, 128);
+BM_FuncWithInput2DimsGPU(colReduction, 6272, 256);
+BM_FuncWithInput2DimsGPU(colReduction, 1568, 204);
+BM_FuncWithInput2DimsGPU(colReduction, 1568, 512);
+BM_FuncWithInput2DimsGPU(colReduction, 1568, 1024);
+BM_FuncWithInput2DimsGPU(colReduction, 1568, 2048);
+BM_FuncWithInput2DimsGPU(fullReduction, 1001, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 2050048, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 2097152, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 2048, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 262144, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 256, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 589824, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 1024, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 524288, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 512, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 2359296, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 1048576, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 131072, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 16384, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 9408, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 64, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 4096, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 36864, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 32768, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 128, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 147456, 1);
+BM_FuncWithInput2DimsGPU(fullReduction, 65536, 1);
+#define BM_FuncGPU(FUNC)                                       \
+  static void BM_##FUNC(int iters, int N) {                    \
+    StopBenchmarkTiming();                                     \
+    Eigen::SyclDevice device(&queue);                          \
+    BenchmarkSuite<Eigen::SyclDevice, float> suite(device, N); \
+    suite.FUNC(iters);                                         \
+  }                                                            \
   BENCHMARK_RANGE(BM_##FUNC, 10, 5000);
 
+BM_FuncGPU(rowReduction);
+BM_FuncGPU(colReduction);
+BM_FuncGPU(fullReduction);
+
 BM_FuncGPU(memcpy);
 BM_FuncGPU(typeCasting);
+BM_FuncGPU(random);
 BM_FuncGPU(slicing);
 BM_FuncGPU(rowChip);
 BM_FuncGPU(colChip);
@@ -28,40 +85,50 @@ BM_FuncGPU(broadcasting);
 BM_FuncGPU(coeffWiseOp);
 BM_FuncGPU(algebraicFunc);
 BM_FuncGPU(transcendentalFunc);
-BM_FuncGPU(rowReduction);
-BM_FuncGPU(colReduction);
-BM_FuncGPU(fullReduction);
-
-
 // Contractions
-#define BM_FuncWithInputDimsGPU(FUNC, D1, D2, D3)                              \
-  static void BM_##FUNC##_##D1##x##D2##x##D3(int iters, int N) {               \
-    StopBenchmarkTiming();                                                     \
-    cl::sycl::gpu_selector selector;                                           \
-    Eigen::QueueInterface queue(selector);                                     \
-    Eigen::SyclDevice device(&queue);                                          \
-    BenchmarkSuite<Eigen::SyclDevice, float> suite(device, D1, D2, D3);        \
-    suite.FUNC(iters);                                                         \
-  }                                                                            \
+#define BM_FuncWithInputDimsGPU(FUNC, D1, D2, D3)                       \
+  static void BM_##FUNC##_##D1##x##D2##x##D3(int iters, int N) {        \
+    StopBenchmarkTiming();                                              \
+    Eigen::SyclDevice device(&queue);                                   \
+    BenchmarkSuite<Eigen::SyclDevice, float> suite(device, D1, D2, D3); \
+    suite.FUNC(iters);                                                  \
+  }                                                                     \
   BENCHMARK_RANGE(BM_##FUNC##_##D1##x##D2##x##D3, 10, 5000);
 
-
 BM_FuncWithInputDimsGPU(contraction, N, N, N);
 BM_FuncWithInputDimsGPU(contraction, 64, N, N);
 BM_FuncWithInputDimsGPU(contraction, N, 64, N);
 BM_FuncWithInputDimsGPU(contraction, N, N, 64);
 
+BM_FuncWithInputDimsGPU(contractionRowMajor, N, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajor, 64, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajor, N, 64, N);
+BM_FuncWithInputDimsGPU(contractionRowMajor, N, N, 64);
+
+BM_FuncWithInputDimsGPU(contractionRowMajorAT, N, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorAT, 64, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorAT, N, 64, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorAT, N, N, 64);
+
+BM_FuncWithInputDimsGPU(contractionRowMajorBT, N, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorBT, 64, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorBT, N, 64, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorBT, N, N, 64);
+
+
+BM_FuncWithInputDimsGPU(contractionRowMajorABT, N, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorABT, 64, N, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorABT, N, 64, N);
+BM_FuncWithInputDimsGPU(contractionRowMajorABT, N, N, 64);
 
 // Convolutions
-#define BM_FuncWithKernelDimsGPU(FUNC, DIM1, DIM2)                             \
-  static void BM_##FUNC##_##DIM1##x##DIM2(int iters, int N) {                  \
-    StopBenchmarkTiming();                                                     \
-    cl::sycl::gpu_selector selector;                                           \
-    Eigen::QueueInterface queue(selector);                                     \
-    Eigen::SyclDevice device(&queue);                                          \
-    BenchmarkSuite<Eigen::SyclDevice, float> suite(device, N);                 \
-    suite.FUNC(iters, DIM1, DIM2);                                             \
-  }                                                                            \
+#define BM_FuncWithKernelDimsGPU(FUNC, DIM1, DIM2)             \
+  static void BM_##FUNC##_##DIM1##x##DIM2(int iters, int N) {  \
+    StopBenchmarkTiming();                                     \
+    Eigen::SyclDevice device(&queue);                          \
+    BenchmarkSuite<Eigen::SyclDevice, float> suite(device, N); \
+    suite.FUNC(iters, DIM1, DIM2);                             \
+  }                                                            \
   BENCHMARK_RANGE(BM_##FUNC##_##DIM1##x##DIM2, 128, 5000);
 
 BM_FuncWithKernelDimsGPU(convolution, 7, 1);