From 524fa4c46f8c5c1bc01f7754188e6883f669a543 Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Fri, 14 Oct 2016 12:09:55 +0100 Subject: Reducing the code by generalising sycl backend functions/structs. --- unsupported/test/cxx11_tensor_sycl.cpp | 92 +++++++++++++++++----------------- 1 file changed, 47 insertions(+), 45 deletions(-) (limited to 'unsupported/test/cxx11_tensor_sycl.cpp') diff --git a/unsupported/test/cxx11_tensor_sycl.cpp b/unsupported/test/cxx11_tensor_sycl.cpp index a4dc9f9fa..0f66cd8f0 100644 --- a/unsupported/test/cxx11_tensor_sycl.cpp +++ b/unsupported/test/cxx11_tensor_sycl.cpp @@ -6,11 +6,13 @@ // Ralph Potter Codeplay Software Ltd. // Luke Iwanski Codeplay Software Ltd. // Contact: +// Benoit Steiner // // This Source Code Form is subject to the terms of the Mozilla // Public License v. 2.0. If a copy of the MPL was not distributed // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. + #define EIGEN_TEST_NO_LONGDOUBLE #define EIGEN_TEST_NO_COMPLEX #define EIGEN_TEST_FUNC cxx11_tensor_sycl @@ -30,7 +32,7 @@ using TestTensor = Tensor; using TestTensorMap = TensorMap>; void test_sycl_cpu() { - cl::sycl::gpu_selector s; + cl::sycl::gpu_selector s; cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { for (const auto& e : l) { try { @@ -40,7 +42,7 @@ void test_sycl_cpu() { } } }); - SyclDevice sycl_device(q); + SyclDevice sycl_device(q); int sizeDim1 = 100; int sizeDim2 = 100; @@ -53,14 +55,14 @@ void test_sycl_cpu() { in1 = in1.random(); in2 = in2.random(); in3 = in3.random(); - TestTensorMap gpu_in1(in1.data(), tensorRange); - TestTensorMap gpu_in2(in2.data(), tensorRange); - TestTensorMap gpu_in3(in3.data(), tensorRange); - TestTensorMap gpu_out(out.data(), tensorRange); - - /// a=1.2f - gpu_in1.device(sycl_device) = gpu_in1.constant(1.2f); - sycl_device.deallocate(in1.data()); + TestTensorMap gpu_in1(in1.data(), tensorRange); + TestTensorMap gpu_in2(in2.data(), tensorRange); + TestTensorMap gpu_in3(in3.data(), tensorRange); + TestTensorMap gpu_out(out.data(), tensorRange); + + /// a=1.2f + gpu_in1.device(sycl_device) = gpu_in1.constant(1.2f); + sycl_device.deallocate(in1.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -68,11 +70,11 @@ void test_sycl_cpu() { } } } - printf("a=1.2f Test passed\n"); + printf("a=1.2f Test passed\n"); - /// a=b*1.2f - gpu_out.device(sycl_device) = gpu_in1 * 1.2f; - sycl_device.deallocate(out.data()); + /// a=b*1.2f + gpu_out.device(sycl_device) = gpu_in1 * 1.2f; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -81,11 +83,11 @@ void test_sycl_cpu() { } } } - printf("a=b*1.2f Test Passed\n"); + printf("a=b*1.2f Test Passed\n"); - /// c=a*b - gpu_out.device(sycl_device) = gpu_in1 * gpu_in2; - sycl_device.deallocate(out.data()); + /// c=a*b + gpu_out.device(sycl_device) = gpu_in1 * gpu_in2; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -95,11 +97,11 @@ void test_sycl_cpu() { } } } - printf("c=a*b Test Passed\n"); + printf("c=a*b Test Passed\n"); - /// c=a+b - gpu_out.device(sycl_device) = gpu_in1 + gpu_in2; - sycl_device.deallocate(out.data()); + /// c=a+b + gpu_out.device(sycl_device) = gpu_in1 + gpu_in2; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -109,11 +111,11 @@ void test_sycl_cpu() { } } } - printf("c=a+b Test Passed\n"); + printf("c=a+b Test Passed\n"); - /// c=a*a - gpu_out.device(sycl_device) = gpu_in1 * gpu_in1; - sycl_device.deallocate(out.data()); + /// c=a*a + gpu_out.device(sycl_device) = gpu_in1 * gpu_in1; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -124,11 +126,11 @@ void test_sycl_cpu() { } } - printf("c= a*a Test Passed\n"); + printf("c= a*a Test Passed\n"); - //a*3.14f + b*2.7f - gpu_out.device(sycl_device) = gpu_in1 * gpu_in1.constant(3.14f) + gpu_in2 * gpu_in2.constant(2.7f); - sycl_device.deallocate(out.data()); + //a*3.14f + b*2.7f + gpu_out.device(sycl_device) = gpu_in1 * gpu_in1.constant(3.14f) + gpu_in2 * gpu_in2.constant(2.7f); + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -138,21 +140,21 @@ void test_sycl_cpu() { } } } - printf("a*3.14f + b*2.7f Test Passed\n"); - - ///d= (a>0.5? b:c) - gpu_out.device(sycl_device) =(gpu_in1 > gpu_in1.constant(0.5f)).select(gpu_in2, gpu_in3); - sycl_device.deallocate(out.data()); - for (int i = 0; i < sizeDim1; ++i) { - for (int j = 0; j < sizeDim2; ++j) { - for (int k = 0; k < sizeDim3; ++k) { - VERIFY_IS_APPROX(out(i, j, k), (in1(i, j, k) > 0.5f) - ? in2(i, j, k) - : in3(i, j, k)); - } - } - } - printf("d= (a>0.5? b:c) Test Passed\n"); + printf("a*3.14f + b*2.7f Test Passed\n"); + + ///d= (a>0.5? b:c) + gpu_out.device(sycl_device) =(gpu_in1 > gpu_in1.constant(0.5f)).select(gpu_in2, gpu_in3); + sycl_device.deallocate(out.data()); + for (int i = 0; i < sizeDim1; ++i) { + for (int j = 0; j < sizeDim2; ++j) { + for (int k = 0; k < sizeDim3; ++k) { + VERIFY_IS_APPROX(out(i, j, k), (in1(i, j, k) > 0.5f) + ? in2(i, j, k) + : in3(i, j, k)); + } + } + } + printf("d= (a>0.5? b:c) Test Passed\n"); } void test_cxx11_tensor_sycl() { -- cgit v1.2.3