From 00f32752f7d0b193c6788691c3cf0b76457a044d Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Thu, 28 Nov 2019 10:08:54 +0000 Subject: [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 --- bench/tensors/tensor_benchmarks_sycl.cc | 133 ++++++++++++++++++++++++-------- 1 file changed, 100 insertions(+), 33 deletions(-) (limited to 'bench/tensors/tensor_benchmarks_sycl.cc') 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 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 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 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 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 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 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 suite(device, N); \ + suite.FUNC(iters, DIM1, DIM2); \ + } \ BENCHMARK_RANGE(BM_##FUNC##_##DIM1##x##DIM2, 128, 5000); BM_FuncWithKernelDimsGPU(convolution, 7, 1); -- cgit v1.2.3