Retain 1.2.1 and 2020 versions of SYCL
This commit is contained in:
Tom Deakin
parent
edcc3e79cd
commit
e077d149dc
@ -9,41 +9,51 @@
|
||||
|
||||
#include <iostream>
|
||||
|
||||
using namespace cl::sycl;
|
||||
|
||||
// Cache list of devices
|
||||
bool cached = false;
|
||||
std::vector<sycl::device> devices;
|
||||
std::vector<device> devices;
|
||||
void getDeviceList(void);
|
||||
|
||||
template <class T>
|
||||
SYCLStream<T>::SYCLStream(const size_t ARRAY_SIZE, const int device_index)
|
||||
: array_size {ARRAY_SIZE},
|
||||
d_a {ARRAY_SIZE},
|
||||
d_b {ARRAY_SIZE},
|
||||
d_c {ARRAY_SIZE},
|
||||
d_sum {1}
|
||||
SYCLStream<T>::SYCLStream(const int ARRAY_SIZE, const int device_index)
|
||||
{
|
||||
if (!cached)
|
||||
getDeviceList();
|
||||
|
||||
array_size = ARRAY_SIZE;
|
||||
|
||||
if (device_index >= devices.size())
|
||||
throw std::runtime_error("Invalid device index");
|
||||
|
||||
sycl::device dev = devices[device_index];
|
||||
|
||||
// Print out device information
|
||||
std::cout << "Using SYCL device " << getDeviceName(device_index) << std::endl;
|
||||
std::cout << "Driver: " << getDeviceDriver(device_index) << std::endl;
|
||||
device dev = devices[device_index];
|
||||
|
||||
// Check device can support FP64 if needed
|
||||
if (sizeof(T) == sizeof(double))
|
||||
{
|
||||
if (!dev.has(sycl::aspect::fp64))
|
||||
{
|
||||
if (dev.get_info<info::device::double_fp_config>().size() == 0) {
|
||||
throw std::runtime_error("Device does not support double precision, please use --float");
|
||||
}
|
||||
}
|
||||
|
||||
queue = std::make_unique<sycl::queue>(dev, sycl::async_handler{[&](sycl::exception_list l)
|
||||
// Determine sensible dot kernel NDRange configuration
|
||||
if (dev.is_cpu())
|
||||
{
|
||||
dot_num_groups = dev.get_info<info::device::max_compute_units>();
|
||||
dot_wgsize = dev.get_info<info::device::native_vector_width_double>() * 2;
|
||||
}
|
||||
else
|
||||
{
|
||||
dot_num_groups = dev.get_info<info::device::max_compute_units>() * 4;
|
||||
dot_wgsize = dev.get_info<info::device::max_work_group_size>();
|
||||
}
|
||||
|
||||
// Print out device information
|
||||
std::cout << "Using SYCL device " << getDeviceName(device_index) << std::endl;
|
||||
std::cout << "Driver: " << getDeviceDriver(device_index) << std::endl;
|
||||
std::cout << "Reduction kernel config: " << dot_num_groups << " groups of size " << dot_wgsize << std::endl;
|
||||
|
||||
queue = new cl::sycl::queue(dev, cl::sycl::async_handler{[&](cl::sycl::exception_list l)
|
||||
{
|
||||
bool error = false;
|
||||
for(auto e: l)
|
||||
@ -52,7 +62,7 @@ SYCLStream<T>::SYCLStream(const size_t ARRAY_SIZE, const int device_index)
|
||||
{
|
||||
std::rethrow_exception(e);
|
||||
}
|
||||
catch (sycl::exception e)
|
||||
catch (cl::sycl::exception e)
|
||||
{
|
||||
std::cout << e.what();
|
||||
error = true;
|
||||
@ -64,22 +74,32 @@ SYCLStream<T>::SYCLStream(const size_t ARRAY_SIZE, const int device_index)
|
||||
}
|
||||
}});
|
||||
|
||||
// No longer need list of devices
|
||||
devices.clear();
|
||||
cached = true;
|
||||
|
||||
|
||||
// Create buffers
|
||||
d_a = new buffer<T>(array_size);
|
||||
d_b = new buffer<T>(array_size);
|
||||
d_c = new buffer<T>(array_size);
|
||||
d_sum = new buffer<T>(dot_num_groups);
|
||||
}
|
||||
|
||||
template <class T>
|
||||
SYCLStream<T>::~SYCLStream()
|
||||
{
|
||||
delete d_a;
|
||||
delete d_b;
|
||||
delete d_c;
|
||||
delete d_sum;
|
||||
delete queue;
|
||||
devices.clear();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::copy()
|
||||
{
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
queue->submit([&](handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::write_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
auto ka = d_a->template get_access<access::mode::read>(cgh);
|
||||
auto kc = d_c->template get_access<access::mode::write>(cgh);
|
||||
cgh.parallel_for<copy_kernel>(range<1>{array_size}, [=](id<1> idx)
|
||||
{
|
||||
kc[idx] = ka[idx];
|
||||
});
|
||||
@ -91,11 +111,11 @@ template <class T>
|
||||
void SYCLStream<T>::mul()
|
||||
{
|
||||
const T scalar = startScalar;
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
queue->submit([&](handler &cgh)
|
||||
{
|
||||
sycl::accessor kb {d_b, cgh, sycl::write_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::read_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
auto kb = d_b->template get_access<access::mode::write>(cgh);
|
||||
auto kc = d_c->template get_access<access::mode::read>(cgh);
|
||||
cgh.parallel_for<mul_kernel>(range<1>{array_size}, [=](id<1> idx)
|
||||
{
|
||||
kb[idx] = scalar * kc[idx];
|
||||
});
|
||||
@ -106,12 +126,12 @@ void SYCLStream<T>::mul()
|
||||
template <class T>
|
||||
void SYCLStream<T>::add()
|
||||
{
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
queue->submit([&](handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::read_only};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::write_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
auto ka = d_a->template get_access<access::mode::read>(cgh);
|
||||
auto kb = d_b->template get_access<access::mode::read>(cgh);
|
||||
auto kc = d_c->template get_access<access::mode::write>(cgh);
|
||||
cgh.parallel_for<add_kernel>(range<1>{array_size}, [=](id<1> idx)
|
||||
{
|
||||
kc[idx] = ka[idx] + kb[idx];
|
||||
});
|
||||
@ -123,12 +143,12 @@ template <class T>
|
||||
void SYCLStream<T>::triad()
|
||||
{
|
||||
const T scalar = startScalar;
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
queue->submit([&](handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::write_only};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::read_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
auto ka = d_a->template get_access<access::mode::write>(cgh);
|
||||
auto kb = d_b->template get_access<access::mode::read>(cgh);
|
||||
auto kc = d_c->template get_access<access::mode::read>(cgh);
|
||||
cgh.parallel_for<triad_kernel>(range<1>{array_size}, [=](id<1> idx)
|
||||
{
|
||||
ka[idx] = kb[idx] + scalar * kc[idx];
|
||||
});
|
||||
@ -140,13 +160,12 @@ template <class T>
|
||||
void SYCLStream<T>::nstream()
|
||||
{
|
||||
const T scalar = startScalar;
|
||||
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
queue->submit([&](handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::read_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
auto ka = d_a->template get_access<access::mode::read_write>(cgh);
|
||||
auto kb = d_b->template get_access<access::mode::read>(cgh);
|
||||
auto kc = d_c->template get_access<access::mode::read>(cgh);
|
||||
cgh.parallel_for<nstream_kernel>(range<1>{array_size}, [=](id<1> idx)
|
||||
{
|
||||
ka[idx] += kb[idx] + scalar * kc[idx];
|
||||
});
|
||||
@ -157,55 +176,73 @@ void SYCLStream<T>::nstream()
|
||||
template <class T>
|
||||
T SYCLStream<T>::dot()
|
||||
{
|
||||
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
queue->submit([&](handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::read_only};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
auto ka = d_a->template get_access<access::mode::read>(cgh);
|
||||
auto kb = d_b->template get_access<access::mode::read>(cgh);
|
||||
auto ksum = d_sum->template get_access<access::mode::write>(cgh);
|
||||
|
||||
cgh.parallel_for(sycl::range<1>{array_size},
|
||||
// Reduction object, to perform summation - initialises the result to zero
|
||||
sycl::reduction(d_sum, cgh, std::plus<T>(), sycl::property::reduction::initialize_to_identity),
|
||||
[=](sycl::id<1> idx, auto& sum)
|
||||
auto wg_sum = accessor<T, 1, access::mode::read_write, access::target::local>(range<1>(dot_wgsize), cgh);
|
||||
|
||||
size_t N = array_size;
|
||||
cgh.parallel_for<dot_kernel>(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
|
||||
{
|
||||
sum += ka[idx] * kb[idx];
|
||||
size_t i = item.get_global_id(0);
|
||||
size_t li = item.get_local_id(0);
|
||||
size_t global_size = item.get_global_range()[0];
|
||||
|
||||
wg_sum[li] = 0.0;
|
||||
for (; i < N; i += global_size)
|
||||
wg_sum[li] += ka[i] * kb[i];
|
||||
|
||||
size_t local_size = item.get_local_range()[0];
|
||||
for (int offset = local_size / 2; offset > 0; offset /= 2)
|
||||
{
|
||||
item.barrier(cl::sycl::access::fence_space::local_space);
|
||||
if (li < offset)
|
||||
wg_sum[li] += wg_sum[li + offset];
|
||||
}
|
||||
|
||||
if (li == 0)
|
||||
ksum[item.get_group(0)] = wg_sum[0];
|
||||
});
|
||||
});
|
||||
|
||||
});
|
||||
|
||||
// Get access on the host, and return a copy of the data (single number)
|
||||
// This will block until the result is available, so no need to wait on the queue.
|
||||
sycl::host_accessor result {d_sum, sycl::read_only};
|
||||
return result[0];
|
||||
T sum = 0.0;
|
||||
auto h_sum = d_sum->template get_access<access::mode::read>();
|
||||
for (int i = 0; i < dot_num_groups; i++)
|
||||
{
|
||||
sum += h_sum[i];
|
||||
}
|
||||
|
||||
return sum;
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::init_arrays(T initA, T initB, T initC)
|
||||
{
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
queue->submit([&](handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::write_only, sycl::no_init};
|
||||
sycl::accessor kb {d_b, cgh, sycl::write_only, sycl::no_init};
|
||||
sycl::accessor kc {d_c, cgh, sycl::write_only, sycl::no_init};
|
||||
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
auto ka = d_a->template get_access<access::mode::write>(cgh);
|
||||
auto kb = d_b->template get_access<access::mode::write>(cgh);
|
||||
auto kc = d_c->template get_access<access::mode::write>(cgh);
|
||||
cgh.parallel_for<init_kernel>(range<1>{array_size}, [=](item<1> item)
|
||||
{
|
||||
ka[idx] = initA;
|
||||
kb[idx] = initB;
|
||||
kc[idx] = initC;
|
||||
auto id = item.get_id(0);
|
||||
ka[id] = initA;
|
||||
kb[id] = initB;
|
||||
kc[id] = initC;
|
||||
});
|
||||
});
|
||||
|
||||
queue->wait();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c)
|
||||
{
|
||||
sycl::host_accessor _a {d_a, sycl::read_only};
|
||||
sycl::host_accessor _b {d_b, sycl::read_only};
|
||||
sycl::host_accessor _c {d_c, sycl::read_only};
|
||||
auto _a = d_a->template get_access<access::mode::read>();
|
||||
auto _b = d_b->template get_access<access::mode::read>();
|
||||
auto _c = d_c->template get_access<access::mode::read>();
|
||||
for (int i = 0; i < array_size; i++)
|
||||
{
|
||||
a[i] = _a[i];
|
||||
@ -217,7 +254,7 @@ void SYCLStream<T>::read_arrays(std::vector<T>& a, std::vector<T>& b, std::vecto
|
||||
void getDeviceList(void)
|
||||
{
|
||||
// Ask SYCL runtime for all devices in system
|
||||
devices = sycl::device::get_devices();
|
||||
devices = cl::sycl::device::get_devices();
|
||||
cached = true;
|
||||
}
|
||||
|
||||
@ -251,7 +288,7 @@ std::string getDeviceName(const int device)
|
||||
|
||||
if (device < devices.size())
|
||||
{
|
||||
name = devices[device].get_info<sycl::info::device::name>();
|
||||
name = devices[device].get_info<info::device::name>();
|
||||
}
|
||||
else
|
||||
{
|
||||
@ -270,7 +307,7 @@ std::string getDeviceDriver(const int device)
|
||||
|
||||
if (device < devices.size())
|
||||
{
|
||||
driver = devices[device].get_info<sycl::info::device::driver_version>();
|
||||
driver = devices[device].get_info<info::device::driver_version>();
|
||||
}
|
||||
else
|
||||
{
|
||||
@ -280,5 +317,6 @@ std::string getDeviceDriver(const int device)
|
||||
return driver;
|
||||
}
|
||||
|
||||
// TODO: Fix kernel names to allow multiple template specializations
|
||||
template class SYCLStream<float>;
|
||||
template class SYCLStream<double>;
|
||||
|
||||
@ -8,13 +8,22 @@
|
||||
#pragma once
|
||||
|
||||
#include <sstream>
|
||||
#include <memory>
|
||||
|
||||
#include "Stream.h"
|
||||
#include "CL/sycl.hpp"
|
||||
|
||||
#include <sycl/sycl.hpp>
|
||||
#define IMPLEMENTATION_STRING "SYCL"
|
||||
|
||||
#define IMPLEMENTATION_STRING "SYCL 2020"
|
||||
namespace sycl_kernels
|
||||
{
|
||||
template <class T> class init;
|
||||
template <class T> class copy;
|
||||
template <class T> class mul;
|
||||
template <class T> class add;
|
||||
template <class T> class triad;
|
||||
template <class T> class nstream;
|
||||
template <class T> class dot;
|
||||
}
|
||||
|
||||
template <class T>
|
||||
class SYCLStream : public Stream<T>
|
||||
@ -24,19 +33,29 @@ class SYCLStream : public Stream<T>
|
||||
size_t array_size;
|
||||
|
||||
// SYCL objects
|
||||
// Queue is a pointer because we allow device selection
|
||||
std::unique_ptr<sycl::queue> queue;
|
||||
cl::sycl::queue *queue;
|
||||
cl::sycl::buffer<T> *d_a;
|
||||
cl::sycl::buffer<T> *d_b;
|
||||
cl::sycl::buffer<T> *d_c;
|
||||
cl::sycl::buffer<T> *d_sum;
|
||||
|
||||
// Buffers
|
||||
sycl::buffer<T> d_a;
|
||||
sycl::buffer<T> d_b;
|
||||
sycl::buffer<T> d_c;
|
||||
sycl::buffer<T> d_sum;
|
||||
// SYCL kernel names
|
||||
typedef sycl_kernels::init<T> init_kernel;
|
||||
typedef sycl_kernels::copy<T> copy_kernel;
|
||||
typedef sycl_kernels::mul<T> mul_kernel;
|
||||
typedef sycl_kernels::add<T> add_kernel;
|
||||
typedef sycl_kernels::triad<T> triad_kernel;
|
||||
typedef sycl_kernels::nstream<T> nstream_kernel;
|
||||
typedef sycl_kernels::dot<T> dot_kernel;
|
||||
|
||||
// NDRange configuration for the dot kernel
|
||||
size_t dot_num_groups;
|
||||
size_t dot_wgsize;
|
||||
|
||||
public:
|
||||
|
||||
SYCLStream(const size_t, const int);
|
||||
~SYCLStream() = default;
|
||||
SYCLStream(const int, const int);
|
||||
~SYCLStream();
|
||||
|
||||
virtual void copy() override;
|
||||
virtual void add() override;
|
||||
|
||||
284
src/sycl/SYCLStream2020.cpp
Normal file
284
src/sycl/SYCLStream2020.cpp
Normal file
@ -0,0 +1,284 @@
|
||||
|
||||
// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith,
|
||||
// University of Bristol HPC
|
||||
//
|
||||
// For full license terms please see the LICENSE file distributed with this
|
||||
// source code
|
||||
|
||||
#include "SYCLStream.h"
|
||||
|
||||
#include <iostream>
|
||||
|
||||
// Cache list of devices
|
||||
bool cached = false;
|
||||
std::vector<sycl::device> devices;
|
||||
void getDeviceList(void);
|
||||
|
||||
template <class T>
|
||||
SYCLStream<T>::SYCLStream(const size_t ARRAY_SIZE, const int device_index)
|
||||
: array_size {ARRAY_SIZE},
|
||||
d_a {ARRAY_SIZE},
|
||||
d_b {ARRAY_SIZE},
|
||||
d_c {ARRAY_SIZE},
|
||||
d_sum {1}
|
||||
{
|
||||
if (!cached)
|
||||
getDeviceList();
|
||||
|
||||
if (device_index >= devices.size())
|
||||
throw std::runtime_error("Invalid device index");
|
||||
|
||||
sycl::device dev = devices[device_index];
|
||||
|
||||
// Print out device information
|
||||
std::cout << "Using SYCL device " << getDeviceName(device_index) << std::endl;
|
||||
std::cout << "Driver: " << getDeviceDriver(device_index) << std::endl;
|
||||
|
||||
// Check device can support FP64 if needed
|
||||
if (sizeof(T) == sizeof(double))
|
||||
{
|
||||
if (!dev.has(sycl::aspect::fp64))
|
||||
{
|
||||
throw std::runtime_error("Device does not support double precision, please use --float");
|
||||
}
|
||||
}
|
||||
|
||||
queue = std::make_unique<sycl::queue>(dev, sycl::async_handler{[&](sycl::exception_list l)
|
||||
{
|
||||
bool error = false;
|
||||
for(auto e: l)
|
||||
{
|
||||
try
|
||||
{
|
||||
std::rethrow_exception(e);
|
||||
}
|
||||
catch (sycl::exception e)
|
||||
{
|
||||
std::cout << e.what();
|
||||
error = true;
|
||||
}
|
||||
}
|
||||
if(error)
|
||||
{
|
||||
throw std::runtime_error("SYCL errors detected");
|
||||
}
|
||||
}});
|
||||
|
||||
// No longer need list of devices
|
||||
devices.clear();
|
||||
cached = true;
|
||||
|
||||
|
||||
}
|
||||
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::copy()
|
||||
{
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::write_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
{
|
||||
kc[idx] = ka[idx];
|
||||
});
|
||||
});
|
||||
queue->wait();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::mul()
|
||||
{
|
||||
const T scalar = startScalar;
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
{
|
||||
sycl::accessor kb {d_b, cgh, sycl::write_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::read_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
{
|
||||
kb[idx] = scalar * kc[idx];
|
||||
});
|
||||
});
|
||||
queue->wait();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::add()
|
||||
{
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::read_only};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::write_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
{
|
||||
kc[idx] = ka[idx] + kb[idx];
|
||||
});
|
||||
});
|
||||
queue->wait();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::triad()
|
||||
{
|
||||
const T scalar = startScalar;
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::write_only};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::read_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
{
|
||||
ka[idx] = kb[idx] + scalar * kc[idx];
|
||||
});
|
||||
});
|
||||
queue->wait();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::nstream()
|
||||
{
|
||||
const T scalar = startScalar;
|
||||
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
sycl::accessor kc {d_c, cgh, sycl::read_only};
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
{
|
||||
ka[idx] += kb[idx] + scalar * kc[idx];
|
||||
});
|
||||
});
|
||||
queue->wait();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
T SYCLStream<T>::dot()
|
||||
{
|
||||
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::read_only};
|
||||
sycl::accessor kb {d_b, cgh, sycl::read_only};
|
||||
|
||||
cgh.parallel_for(sycl::range<1>{array_size},
|
||||
// Reduction object, to perform summation - initialises the result to zero
|
||||
sycl::reduction(d_sum, cgh, std::plus<T>(), sycl::property::reduction::initialize_to_identity),
|
||||
[=](sycl::id<1> idx, auto& sum)
|
||||
{
|
||||
sum += ka[idx] * kb[idx];
|
||||
});
|
||||
|
||||
});
|
||||
|
||||
// Get access on the host, and return a copy of the data (single number)
|
||||
// This will block until the result is available, so no need to wait on the queue.
|
||||
sycl::host_accessor result {d_sum, sycl::read_only};
|
||||
return result[0];
|
||||
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::init_arrays(T initA, T initB, T initC)
|
||||
{
|
||||
queue->submit([&](sycl::handler &cgh)
|
||||
{
|
||||
sycl::accessor ka {d_a, cgh, sycl::write_only, sycl::no_init};
|
||||
sycl::accessor kb {d_b, cgh, sycl::write_only, sycl::no_init};
|
||||
sycl::accessor kc {d_c, cgh, sycl::write_only, sycl::no_init};
|
||||
|
||||
cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx)
|
||||
{
|
||||
ka[idx] = initA;
|
||||
kb[idx] = initB;
|
||||
kc[idx] = initC;
|
||||
});
|
||||
});
|
||||
|
||||
queue->wait();
|
||||
}
|
||||
|
||||
template <class T>
|
||||
void SYCLStream<T>::read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c)
|
||||
{
|
||||
sycl::host_accessor _a {d_a, sycl::read_only};
|
||||
sycl::host_accessor _b {d_b, sycl::read_only};
|
||||
sycl::host_accessor _c {d_c, sycl::read_only};
|
||||
for (int i = 0; i < array_size; i++)
|
||||
{
|
||||
a[i] = _a[i];
|
||||
b[i] = _b[i];
|
||||
c[i] = _c[i];
|
||||
}
|
||||
}
|
||||
|
||||
void getDeviceList(void)
|
||||
{
|
||||
// Ask SYCL runtime for all devices in system
|
||||
devices = sycl::device::get_devices();
|
||||
cached = true;
|
||||
}
|
||||
|
||||
void listDevices(void)
|
||||
{
|
||||
getDeviceList();
|
||||
|
||||
// Print device names
|
||||
if (devices.size() == 0)
|
||||
{
|
||||
std::cerr << "No devices found." << std::endl;
|
||||
}
|
||||
else
|
||||
{
|
||||
std::cout << std::endl;
|
||||
std::cout << "Devices:" << std::endl;
|
||||
for (int i = 0; i < devices.size(); i++)
|
||||
{
|
||||
std::cout << i << ": " << getDeviceName(i) << std::endl;
|
||||
}
|
||||
std::cout << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
std::string getDeviceName(const int device)
|
||||
{
|
||||
if (!cached)
|
||||
getDeviceList();
|
||||
|
||||
std::string name;
|
||||
|
||||
if (device < devices.size())
|
||||
{
|
||||
name = devices[device].get_info<sycl::info::device::name>();
|
||||
}
|
||||
else
|
||||
{
|
||||
throw std::runtime_error("Error asking for name for non-existant device");
|
||||
}
|
||||
|
||||
return name;
|
||||
}
|
||||
|
||||
std::string getDeviceDriver(const int device)
|
||||
{
|
||||
if (!cached)
|
||||
getDeviceList();
|
||||
|
||||
std::string driver;
|
||||
|
||||
if (device < devices.size())
|
||||
{
|
||||
driver = devices[device].get_info<sycl::info::device::driver_version>();
|
||||
}
|
||||
else
|
||||
{
|
||||
throw std::runtime_error("Error asking for driver for non-existant device");
|
||||
}
|
||||
|
||||
return driver;
|
||||
}
|
||||
|
||||
template class SYCLStream<float>;
|
||||
template class SYCLStream<double>;
|
||||
54
src/sycl/SYCLStream2020.h
Normal file
54
src/sycl/SYCLStream2020.h
Normal file
@ -0,0 +1,54 @@
|
||||
|
||||
// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith,
|
||||
// University of Bristol HPC
|
||||
//
|
||||
// For full license terms please see the LICENSE file distributed with this
|
||||
// source code
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <sstream>
|
||||
#include <memory>
|
||||
|
||||
#include "Stream.h"
|
||||
|
||||
#include <sycl/sycl.hpp>
|
||||
|
||||
#define IMPLEMENTATION_STRING "SYCL 2020"
|
||||
|
||||
template <class T>
|
||||
class SYCLStream : public Stream<T>
|
||||
{
|
||||
protected:
|
||||
// Size of arrays
|
||||
size_t array_size;
|
||||
|
||||
// SYCL objects
|
||||
// Queue is a pointer because we allow device selection
|
||||
std::unique_ptr<sycl::queue> queue;
|
||||
|
||||
// Buffers
|
||||
sycl::buffer<T> d_a;
|
||||
sycl::buffer<T> d_b;
|
||||
sycl::buffer<T> d_c;
|
||||
sycl::buffer<T> d_sum;
|
||||
|
||||
public:
|
||||
|
||||
SYCLStream(const size_t, const int);
|
||||
~SYCLStream() = default;
|
||||
|
||||
virtual void copy() override;
|
||||
virtual void add() override;
|
||||
virtual void mul() override;
|
||||
virtual void triad() override;
|
||||
virtual void nstream() override;
|
||||
virtual T dot() override;
|
||||
|
||||
virtual void init_arrays(T initA, T initB, T initC) override;
|
||||
virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
|
||||
|
||||
};
|
||||
|
||||
// Populate the devices list
|
||||
void getDeviceList(void);
|
||||
Loading…
Reference in New Issue
Block a user