devtools/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2026-04-10 11:34:33 +08:00
Code Issues Packages Projects Releases Wiki Activity

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.

Browse Source
This commit is contained in:
Mehdi Goli
2016-11-08 17:08:02 +00:00
parent dad177be01
commit d57430dd73
11 changed files with 328 additions and 341 deletions
Download Patch File Download Diff File Expand all files Collapse all files

85
unsupported/test/cxx11_tensor_broadcast_sycl.cpp
View File

@@ -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;
// 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);
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);
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)));
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));
}
}
}
});
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];
Tensor<float, in_range.size()> input(in_range);
Tensor<float, out_range.size()> output(out_range);
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 (size_t i = 0; i < in_range.size(); ++i)
VERIFY_IS_EQUAL(output.dimension(i), out_range[i]);
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");
}
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));
}