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

Merge branch 'refactor' of https://github.com/UoB-hpc/gpu-stream into refactor

Browse Source
This commit is contained in:
Matthew Martineau 2016-05-06 10:54:18 +01:00
commit 57189e7ca5
4 changed files with 46 additions and 30 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
CMakeLists.txt
View File

@ -40,9 +40,18 @@ endif ()
# TODO: Find SYCL implementations somehow
if (true)
add_executable(gpu-stream-sycl main.cpp SYCLStream.cpp)
target_compile_definitions(gpu-stream-sycl PUBLIC SYCL)
set_property(TARGET gpu-stream-sycl PROPERTY CXX_STANDARD 14)
# ComputeCpp
# TODO: Sort this out properly!
add_custom_target(gpu-stream-sycl
COMMAND compute++ ${CMAKE_CURRENT_SOURCE_DIR}/SYCLStream.cpp -sycl -O2 -emit-llvm -o SYCLStream.bc -c
COMMAND /usr/bin/c++ ${CMAKE_CURRENT_SOURCE_DIR}/SYCLStream.cpp -O2 -std=c++11 -include SYCLStream.sycl -o SYCLStream.o -c
COMMAND /usr/bin/c++ ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp -O2 -std=c++11 SYCLStream.o -include SYCLStream.sycl -lSYCL -lOpenCL -o gpu-stream-sycl -DSYCL
)
# triSYCL
#add_executable(gpu-stream-sycl main.cpp SYCLStream.cpp)
#target_compile_definitions(gpu-stream-sycl PUBLIC SYCL)
#set_property(TARGET gpu-stream-sycl PROPERTY CXX_STANDARD 14)
else ()
message("Skipping SYCL...")
endif ()

52
SYCLStream.cpp
View File

@ -17,14 +17,17 @@ SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
array_size = ARRAY_SIZE;
// Create buffers
d_a = buffer<T>(array_size);
d_b = buffer<T>(array_size);
d_c = buffer<T>(array_size);
d_a = new buffer<T>(array_size);
d_b = new buffer<T>(array_size);
d_c = new buffer<T>(array_size);
}
template <class T>
SYCLStream<T>::~SYCLStream()
{
delete d_a;
delete d_b;
delete d_c;
}
template <class T>
@ -32,9 +35,9 @@ void SYCLStream<T>::copy()
{
queue.submit([&](handler &cgh)
{
auto ka = d_a.template get_access<access::read>(cgh);
auto kc = d_c.template get_access<access::write>(cgh);
cgh.parallel_for(range<1>{array_size}, [=](id<1> index)
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<class copy>(range<1>{array_size}, [=](id<1> index)
{
kc[index] = ka[index];
});
@ -48,9 +51,9 @@ void SYCLStream<T>::mul()
const T scalar = 3.0;
queue.submit([&](handler &cgh)
{
auto kb = d_b.template get_access<access::write>(cgh);
auto kc = d_c.template get_access<access::read>(cgh);
cgh.parallel_for(range<1>{array_size}, [=](id<1> index)
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<class mul>(range<1>{array_size}, [=](id<1> index)
{
kb[index] = scalar * kc[index];
});
@ -63,10 +66,10 @@ void SYCLStream<T>::add()
{
queue.submit([&](handler &cgh)
{
auto ka = d_a.template get_access<access::read>(cgh);
auto kb = d_b.template get_access<access::read>(cgh);
auto kc = d_c.template get_access<access::write>(cgh);
cgh.parallel_for(range<1>{array_size}, [=](id<1> index)
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<class add>(range<1>{array_size}, [=](id<1> index)
{
kc[index] = ka[index] + kb[index];
});
@ -80,10 +83,10 @@ void SYCLStream<T>::triad()
const T scalar = 3.0;
queue.submit([&](handler &cgh)
{
auto ka = d_a.template get_access<access::write>(cgh);
auto kb = d_b.template get_access<access::read>(cgh);
auto kc = d_c.template get_access<access::read>(cgh);
cgh.parallel_for(range<1>{array_size}, [=](id<1> index){
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<class triad>(range<1>{array_size}, [=](id<1> index){
ka[index] = kb[index] + scalar * kc[index];
});
});
@ -93,9 +96,9 @@ void SYCLStream<T>::triad()
template <class T>
void SYCLStream<T>::write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
{
auto _a = d_a.template get_access<access::write>();
auto _b = d_b.template get_access<access::write>();
auto _c = d_c.template get_access<access::write>();
auto _a = d_a->template get_access<access::mode::write, access::target::host_buffer>();
auto _b = d_b->template get_access<access::mode::write, access::target::host_buffer>();
auto _c = d_c->template get_access<access::mode::write, access::target::host_buffer>();
for (int i = 0; i < array_size; i++)
{
_a[i] = a[i];
@ -107,9 +110,9 @@ void SYCLStream<T>::write_arrays(const std::vector<T>& a, const std::vector<T>&
template <class T>
void SYCLStream<T>::read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c)
{
auto _a = d_a.template get_access<access::read>();
auto _b = d_b.template get_access<access::read>();
auto _c = d_c.template get_access<access::read>();
auto _a = d_a->template get_access<access::mode::read, access::target::host_buffer>();
auto _b = d_b->template get_access<access::mode::read, access::target::host_buffer>();
auto _c = d_c->template get_access<access::mode::read, access::target::host_buffer>();
for (int i = 0; i < array_size; i++)
{
a[i] = _a[i];
@ -140,5 +143,6 @@ std::string getDeviceDriver(const int device)
}
template class SYCLStream<float>;
// TODO: Fix kernel names to allow multiple template specializations
//template class SYCLStream<float>;
template class SYCLStream<double>;

6
SYCLStream.h
View File

@ -22,9 +22,9 @@ class SYCLStream : public Stream<T>
// SYCL objects
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_a;
cl::sycl::buffer<T> *d_b;
cl::sycl::buffer<T> *d_c;
public:

3
main.cpp
View File

@ -57,9 +57,12 @@ int main(int argc, char *argv[])
parseArguments(argc, argv);
// TODO: Fix SYCL to allow multiple template specializations
#ifndef SYCL
if (use_float)
run<float>();
else
#endif
run<double>();
}