Skip to content

Commit

Permalink
Use sycl2020
Browse files Browse the repository at this point in the history
  • Loading branch information
@sbalint98
sbalint98 committed Feb 20, 2024
1 parent 10defa6 commit 5e23f83
Show file tree
Hide file tree
Showing 54 changed files with 677 additions and 670 deletions.
16 changes: 8 additions & 8 deletions tests/compiler/cbs/accumulator_for.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,31 +6,31 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
for(size_t i = 0; i < global_size; ++i)
{
host_buf.push_back(static_cast<int>(i));
}

{
cl::sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};
auto scratch = sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};

cgh.parallel_for<class dynamic_local_memory_reduction>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept {
const auto lid = item.get_local_id(0);

scratch[lid] = acc[item.get_global_id()];
Expand Down
16 changes: 8 additions & 8 deletions tests/compiler/cbs/add_modulo.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,31 +6,31 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
for(size_t i = 0; i < global_size; ++i)
{
host_buf.push_back(static_cast<int>(i));
}

{
cl::sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};
auto scratch = sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};

cgh.parallel_for<class dynamic_local_memory_reduction>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept {
const auto lid = item.get_local_id(0);
const auto group_size = item.get_local_range(0);

Expand Down
22 changes: 11 additions & 11 deletions tests/compiler/cbs/babelstream_dot.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,15 +7,15 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t dot_wgsize = 4;
constexpr size_t dot_num_groups = 32;
constexpr size_t array_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
std::vector<int> host_buf2;
std::vector<int> host_outbuf(dot_num_groups);
Expand All @@ -26,22 +26,22 @@ int main()
}

{
cl::sycl::buffer<int, 1> d_a{host_buf.data(), host_buf.size()};
cl::sycl::buffer<int, 1> d_b{host_buf2.data(), host_buf2.size()};
cl::sycl::buffer<int, 1> d_sum{host_outbuf.data(), host_outbuf.size()};
sycl::buffer<int, 1> d_a{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> d_b{host_buf2.data(), host_buf2.size()};
sycl::buffer<int, 1> d_sum{host_outbuf.data(), host_outbuf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto ka = d_a.template get_access<mode::read>(cgh);
auto kb = d_b.template get_access<mode::read>(cgh);
auto ksum = d_sum.template get_access<mode::discard_write>(cgh);

auto wg_sum = cl::sycl::accessor<int, 1, mode::read_write, target::local>(cl::sycl::range<1>(dot_wgsize), cgh);
auto wg_sum = sycl::accessor<int, 1, mode::read_write, target::local>(sycl::range<1>(dot_wgsize), cgh);

size_t N = array_size;
cgh.parallel_for<class dot_kernel>(
cl::sycl::nd_range<1>{dot_num_groups * dot_wgsize, dot_wgsize},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{dot_num_groups * dot_wgsize, dot_wgsize},
[=](sycl::nd_item<1> item) noexcept {
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];
Expand All @@ -53,7 +53,7 @@ int main()
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);
item.barrier(sycl::access::fence_space::local_space);
if(li < offset)
wg_sum[li] += wg_sum[li + offset];
}
Expand Down
20 changes: 10 additions & 10 deletions tests/compiler/cbs/cond_between_barriers.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5,14 +5,14 @@

#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<char> host_buf(global_size);

for(size_t i = 0; i < 2 * local_size; ++i)
Expand All @@ -24,28 +24,28 @@ int main()
host_buf[2 * local_size + 10] = false;

{
cl::sycl::buffer<char, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<char, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<bool, 1, mode::read_write, target::local>{1, cgh};
auto scratch = sycl::accessor<bool, 1, mode::read_write, target::local>{1, cgh};

cgh.parallel_for<class test_kernel>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept
{
auto g = item.get_group();
const auto id = item.get_global_id();
const bool pred = acc[id];

scratch[0] = false;
cl::sycl::group_barrier(g);
sycl::group_barrier(g);

if (pred)
scratch[0] = pred;

cl::sycl::group_barrier(g);
sycl::group_barrier(g);

acc[id] = scratch[0];
});
Expand Down
16 changes: 8 additions & 8 deletions tests/compiler/cbs/conds.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,31 +6,31 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
for(size_t i = 0; i < global_size; ++i)
{
host_buf.push_back(static_cast<int>(i));
}

{
cl::sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};
auto scratch = sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};

cgh.parallel_for<class dynamic_local_memory_reduction>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept {
const auto lid = item.get_local_id(0);
const auto groupId = item.get_group(0);

Expand Down
16 changes: 8 additions & 8 deletions tests/compiler/cbs/conds_in_for.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,31 +6,31 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
for(size_t i = 0; i < global_size; ++i)
{
host_buf.push_back(static_cast<int>(i));
}

{
cl::sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};
auto scratch = sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};

cgh.parallel_for<class dynamic_local_memory_reduction>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept {
const auto lid = item.get_local_id(0);
const auto groupId = item.get_group(0);

Expand Down
16 changes: 8 additions & 8 deletions tests/compiler/cbs/const_init_accumulator_for.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,31 +6,31 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
for(size_t i = 0; i < global_size; ++i)
{
host_buf.push_back(static_cast<int>(i));
}

{
cl::sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};
auto scratch = sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};

cgh.parallel_for<class dynamic_local_memory_reduction>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept {
const auto lid = item.get_local_id(0);

scratch[lid] = acc[item.get_global_id()];
Expand Down
16 changes: 8 additions & 8 deletions tests/compiler/cbs/for_in_cond.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,31 +7,31 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
for(size_t i = 0; i < global_size; ++i)
{
host_buf.push_back(static_cast<int>(i));
}

{
cl::sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};
auto scratch = sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};

cgh.parallel_for<class dynamic_local_memory_reduction>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept {
const auto lid = item.get_local_id(0);
const auto group_size = item.get_local_range(0);

Expand Down
18 changes: 9 additions & 9 deletions tests/compiler/cbs/group_barrier.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,38 +6,38 @@
#include <array>
#include <iostream>

#include <CL/sycl.hpp>
#include <sycl/sycl.hpp>

int main()
{
constexpr size_t local_size = 256;
constexpr size_t global_size = 1024;

cl::sycl::queue queue;
sycl::queue queue;
std::vector<int> host_buf;
for(size_t i = 0; i < global_size; ++i)
{
host_buf.push_back(static_cast<int>(i));
}

{
cl::sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};
sycl::buffer<int, 1> buf{host_buf.data(), host_buf.size()};

queue.submit([&](cl::sycl::handler &cgh) {
using namespace cl::sycl::access;
queue.submit([&](sycl::handler &cgh) {
using namespace sycl::access;
auto acc = buf.get_access<mode::read_write>(cgh);
auto scratch = cl::sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};
auto scratch = sycl::accessor<int, 1, mode::read_write, target::local>{local_size, cgh};

cgh.parallel_for<class dynamic_local_memory_reduction>(
cl::sycl::nd_range<1>{global_size, local_size},
[=](cl::sycl::nd_item<1> item) noexcept {
sycl::nd_range<1>{global_size, local_size},
[=](sycl::nd_item<1> item) noexcept {
const auto lid = item.get_local_id(0);
const auto group_size = item.get_local_range(0);

scratch[lid] = acc[item.get_global_id()];
for(size_t i = group_size / 2; i > 0; i /= 2)
{
cl::sycl::group_barrier(item.get_group());
sycl::group_barrier(item.get_group());
if(lid < i)
scratch[lid] += scratch[lid + i];
}
Expand Down

0 comments on commit 5e23f83

Please sign in to comment.