cory-master/BabelStream
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[SYCL] Actually use device_index to select device

Browse Source
This commit is contained in:
James Price 2016-05-08 21:35:24 +01:00
parent 3b3f6dfc26
commit 6d913bab4b
2 changed files with 20 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

27
SYCLStream.cpp
View File

@ -21,12 +21,21 @@ void getDeviceList(void);
template <class T>
SYCLStream<T>::SYCLStream(const unsigned 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");
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;
queue = new cl::sycl::queue(dev);
// Create buffers
d_a = new buffer<T>(array_size);
d_b = new buffer<T>(array_size);
@ -39,12 +48,14 @@ SYCLStream<T>::~SYCLStream()
delete d_a;
delete d_b;
delete d_c;
delete queue;
}
template <class T>
void SYCLStream<T>::copy()
{
queue.submit([&](handler &cgh)
queue->submit([&](handler &cgh)
{
auto ka = d_a->template get_access<access::mode::read>(cgh);
auto kc = d_c->template get_access<access::mode::write>(cgh);
@ -53,14 +64,14 @@ void SYCLStream<T>::copy()
kc[item.get_global()] = ka[item.get_global()];
});
});
queue.wait();
queue->wait();
}
template <class T>
void SYCLStream<T>::mul()
{
const T scalar = 3.0;
queue.submit([&](handler &cgh)
queue->submit([&](handler &cgh)
{
auto kb = d_b->template get_access<access::mode::write>(cgh);
auto kc = d_c->template get_access<access::mode::read>(cgh);
@ -69,13 +80,13 @@ void SYCLStream<T>::mul()
kb[item.get_global()] = scalar * kc[item.get_global()];
});
});
queue.wait();
queue->wait();
}
template <class T>
void SYCLStream<T>::add()
{
queue.submit([&](handler &cgh)
queue->submit([&](handler &cgh)
{
auto ka = d_a->template get_access<access::mode::read>(cgh);
auto kb = d_b->template get_access<access::mode::read>(cgh);
@ -85,14 +96,14 @@ void SYCLStream<T>::add()
kc[item.get_global()] = ka[item.get_global()] + kb[item.get_global()];
});
});
queue.wait();
queue->wait();
}
template <class T>
void SYCLStream<T>::triad()
{
const T scalar = 3.0;
queue.submit([&](handler &cgh)
queue->submit([&](handler &cgh)
{
auto ka = d_a->template get_access<access::mode::write>(cgh);
auto kb = d_b->template get_access<access::mode::read>(cgh);
@ -102,7 +113,7 @@ void SYCLStream<T>::triad()
ka[item.get_global()] = kb[item.get_global()] + scalar * kc[item.get_global()];
});
});
queue.wait();
queue->wait();
}
template <class T>

2
SYCLStream.h
View File

@ -21,7 +21,7 @@ class SYCLStream : public Stream<T>
unsigned int array_size;
// SYCL objects
cl::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;