Converting all sycl buffers to uninitialised device only buffers; adding memcpyHostToDevice and memcpyDeviceToHost on syclDevice; modifying all examples to obey the new rules; moving sycl queue creating to the device based on Benoit suggestion; removing the sycl specefic condition for returning m_result in TensorReduction.h according to Benoit suggestion.

1 files changed, 37 insertions, 42 deletions

diff --git a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp
index ecebf7d68..7201bfe37 100644
--- a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp
+++ b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp
@@ -25,55 +25,50 @@ using Eigen::SyclDevice;
 using Eigen::Tensor;
 using Eigen::TensorMap;
 
-// Types used in tests:
-using TestTensor = Tensor<float, 3>;
-using TestTensorMap = TensorMap<Tensor<float, 3>>;
-static void test_broadcast_sycl(){
+static void test_broadcast_sycl(const Eigen::SyclDevice &sycl_device){
 
-	cl::sycl::gpu_selector s;
-  cl::sycl::queue q(s, [=](cl::sycl::exception_list l) {
-    for (const auto& e : l) {
-      try {
-        std::rethrow_exception(e);
-      } catch (cl::sycl::exception e) {
-        std::cout << e.what() << std::endl;
-      }
-    }
-  });
-	SyclDevice sycl_device(q);
-	// BROADCAST test:
-	array<int, 4> in_range   = {{2, 3, 5, 7}};
-	array<int, in_range.size()> broadcasts = {{2, 3, 1, 4}};
-	array<int, in_range.size()> out_range;  // = in_range * broadcasts
-	for (size_t i = 0; i < out_range.size(); ++i)
-		out_range[i] = in_range[i] * broadcasts[i];
+  // BROADCAST test:
+  array<int, 4> in_range   = {{2, 3, 5, 7}};
+  array<int, 4> broadcasts = {{2, 3, 1, 4}};
+  array<int, 4> out_range;  // = in_range * broadcasts
+  for (size_t i = 0; i < out_range.size(); ++i)
+    out_range[i] = in_range[i] * broadcasts[i];
+
+  Tensor<float, 4>  input(in_range);
+  Tensor<float, 4> out(out_range);
 
-	Tensor<float, in_range.size()>  input(in_range);
-	Tensor<float, out_range.size()> output(out_range);
+  for (size_t i = 0; i < in_range.size(); ++i)
+    VERIFY_IS_EQUAL(out.dimension(i), out_range[i]);
 
-	for (int i = 0; i < input.size(); ++i)
-		input(i) = static_cast<float>(i);
 
-	TensorMap<decltype(input)>  gpu_in(input.data(), in_range);
-	TensorMap<decltype(output)> gpu_out(output.data(), out_range);
-		gpu_out.device(sycl_device) = gpu_in.broadcast(broadcasts);
-		sycl_device.deallocate(output.data());
+  for (int i = 0; i < input.size(); ++i)
+    input(i) = static_cast<float>(i);
 
-	for (size_t i = 0; i < in_range.size(); ++i)
-		VERIFY_IS_EQUAL(output.dimension(i), out_range[i]);
+  float * gpu_in_data  = static_cast<float*>(sycl_device.allocate(input.dimensions().TotalSize()*sizeof(float)));
+  float * gpu_out_data  = static_cast<float*>(sycl_device.allocate(out.dimensions().TotalSize()*sizeof(float)));
 
-	for (int i = 0; i < 4; ++i) {
-		for (int j = 0; j < 9; ++j) {
-			for (int k = 0; k < 5; ++k) {
-				for (int l = 0; l < 28; ++l) {
-					VERIFY_IS_APPROX(input(i%2,j%3,k%5,l%7), output(i,j,k,l));
-				}
-			}
-		}
-	}
-	printf("Broadcast Test Passed\n");
+  TensorMap<Tensor<float, 4>>  gpu_in(gpu_in_data, in_range);
+  TensorMap<Tensor<float, 4>> gpu_out(gpu_out_data, out_range);
+  sycl_device.memcpyHostToDevice(gpu_in_data, input.data(),(input.dimensions().TotalSize())*sizeof(float));
+  gpu_out.device(sycl_device) = gpu_in.broadcast(broadcasts);
+  sycl_device.memcpyDeviceToHost(out.data(), gpu_out_data,(out.dimensions().TotalSize())*sizeof(float));
+
+  for (int i = 0; i < 4; ++i) {
+    for (int j = 0; j < 9; ++j) {
+      for (int k = 0; k < 5; ++k) {
+        for (int l = 0; l < 28; ++l) {
+          VERIFY_IS_APPROX(input(i%2,j%3,k%5,l%7), out(i,j,k,l));
+        }
+      }
+    }
+  }
+  printf("Broadcast Test Passed\n");
+  sycl_device.deallocate(gpu_in_data);
+  sycl_device.deallocate(gpu_out_data);
 }
 
 void test_cxx11_tensor_broadcast_sycl() {
-  CALL_SUBTEST(test_broadcast_sycl());
+  cl::sycl::gpu_selector s;
+  Eigen::SyclDevice sycl_device(s);
+  CALL_SUBTEST(test_broadcast_sycl(sycl_device));
 }