Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions sycl/test-e2e/AddressCast/khr_static_addrspace_cast.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -24,10 +24,10 @@ int main() {
Queue
.submit([&](sycl::handler &cgh) {
auto GlobalAccessor =
GlobalBuffer.get_access<sycl::access::mode::read_write>(cgh);
GlobalBuffer.get_access<sycl::access_mode::read_write>(cgh);
auto LocalAccessor = sycl::local_accessor<int>(1, cgh);
auto ResultAccessor =
ResultBuffer.get_access<sycl::access::mode::write>(cgh);
ResultBuffer.get_access<sycl::access_mode::write>(cgh);
cgh.parallel_for<class Kernel>(
sycl::nd_range<1>(NItems, 1), [=](sycl::nd_item<1> Item) {
bool Success = true;
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/Assert/assert_in_kernels.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@ void kernelFunc1(int *Buf, int wiID) {

void assertTest1(queue &Q, buffer<int, 1> &Buf) {
Q.submit([&](handler &CGH) {
auto Acc = Buf.template get_access<mode::read_write>(CGH);
auto Acc = Buf.template get_access<access_mode::read_write>(CGH);

CGH.parallel_for<class Kernel_1>(
Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc1(&Acc[0], wiID); });
Expand All @@ -29,7 +29,7 @@ void kernelFunc2(int *Buf, int wiID) {

void assertTest2(queue &Q, buffer<int, 1> &Buf) {
Q.submit([&](handler &CGH) {
auto Acc = Buf.template get_access<mode::read_write>(CGH);
auto Acc = Buf.template get_access<access_mode::read_write>(CGH);

CGH.parallel_for<class Kernel_2>(
Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc2(&Acc[0], wiID); });
Expand All @@ -44,7 +44,7 @@ void kernelFunc3(int *Buf, int wiID) {

void assertTest3(queue &Q, buffer<int, 1> &Buf) {
Q.submit([&](handler &CGH) {
auto Acc = Buf.template get_access<mode::read_write>(CGH);
auto Acc = Buf.template get_access<access_mode::read_write>(CGH);

CGH.parallel_for<class Kernel_3>(
Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc3(&Acc[0], wiID); });
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ template <class kernel_name> void enqueueKernel(queue *Q) {
sycl::buffer<int, 1> Buf(numOfItems);

Q->submit([&](handler &CGH) {
auto Acc = Buf.template get_access<mode::read_write>(CGH);
auto Acc = Buf.template get_access<access_mode::read_write>(CGH);

CGH.parallel_for<kernel_name>(numOfItems, [=](sycl::id<1> wiID) {
Acc[wiID] = 0;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/AtomicRef/assignment.h
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@ void assignment_test(queue q, size_t N) {
buffer<T> assignment_buf(&assignment, 1);
q.submit([&](handler &cgh) {
auto st =
assignment_buf.template get_access<access::mode::read_write>(cgh);
assignment_buf.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for<assignment_kernel<AtomicRef, address_space, T>>(
range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/AtomicRef/compare_exchange.h
Original file line number Diff line number Diff line change
Expand Up @@ -33,10 +33,10 @@ void compare_exchange_local_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto compare_exchange =
compare_exchange_buf.template get_access<access::mode::read_write>(
compare_exchange_buf.template get_access<access_mode::read_write>(
cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
local_accessor<T, 1> loc(1, cgh);
Comment thread
KornevNikita marked this conversation as resolved.

cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
Expand Down Expand Up @@ -88,10 +88,10 @@ void compare_exchange_global_test(queue q, size_t N) {

q.submit([&](handler &cgh) {
auto exc =
compare_exchange_buf.template get_access<access::mode::read_write>(
compare_exchange_buf.template get_access<access_mode::read_write>(
cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/AtomicRef/max.h
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,9 @@ void max_local_test(queue q, size_t N) {
buffer<T> cum_buf(&cum, 1);
buffer<T> output_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
auto cum = cum_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
local_accessor<T, 1> loc(1, cgh);
Comment thread
KornevNikita marked this conversation as resolved.

cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
Expand Down Expand Up @@ -83,9 +83,9 @@ void max_global_test(queue q, size_t N) {
buffer<T> output_buf(output.data(), output.size());

q.submit([&](handler &cgh) {
auto val = val_buf.template get_access<access::mode::read_write>(cgh);
auto val = val_buf.template get_access<access_mode::read_write>(cgh);
auto out =
output_buf.template get_access<access::mode::discard_write>(cgh);
output_buf.template get_access<access_mode::discard_write>(cgh);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
int gid = it.get_id(0);
auto atm = AtomicRef < T,
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/AtomicRef/store.h
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ void store_global_test(queue q, size_t N) {
{
buffer<T> store_buf(&store, 1);
q.submit([&](handler &cgh) {
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
auto st = store_buf.template get_access<access_mode::read_write>(cgh);
cgh.parallel_for(range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
auto atm = AtomicRef<T, memory_order::relaxed, scope, space>(st[0]);
Expand Down Expand Up @@ -82,7 +82,7 @@ void store_local_test(queue q, size_t N) {
{
buffer<T> store_buf(&store, 1);
q.submit([&](handler &cgh) {
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
auto st = store_buf.template get_access<access_mode::read_write>(cgh);
local_accessor<T, 1> loc(1, cgh);
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
size_t gid = it.get_global_id(0);
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/access_to_subset.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@
#include <sycl/detail/core.hpp>

using namespace sycl;
using acc_w = accessor<int, 2, access::mode::write, access::target::device>;
using acc_w = accessor<int, 2, access_mode::write, access::target::device>;

int main() {

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/bit_cast/bit_cast.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
#include <math.h>
#include <type_traits>

constexpr sycl::access::mode sycl_write = sycl::access::mode::write;
constexpr sycl::access_mode sycl_write = sycl::access_mode::write;

template <typename To, typename From> class BitCastKernel;

Expand Down
18 changes: 9 additions & 9 deletions sycl/test-e2e/Basic/buffer/buffer_full_copy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@ void check_copy_device_to_host(sycl::queue &Queue) {

// Submit kernel where you'll fill the buffer
Queue.submit([&](sycl::handler &cgh) {
auto acc = simple_buffer.get_access<sycl::access::mode::write>(cgh);
auto acc = simple_buffer.get_access<sycl::access_mode::write>(cgh);
cgh.fill(acc, 13);
});

Expand Down Expand Up @@ -69,8 +69,8 @@ void check_fill(sycl::queue &Queue) {
}

auto e = Queue.submit([&](sycl::handler &cgh) {
auto a = buf_1.template get_access<sycl::access::mode::write>(cgh, size / 2,
offset);
auto a = buf_1.template get_access<sycl::access_mode::write>(cgh, size / 2,
offset);
cgh.fill(a, (float)1337.0);
});
e.wait();
Expand Down Expand Up @@ -102,8 +102,8 @@ void check_copy_host_to_device(sycl::queue &Queue) {
expected_res_2[i] = expected_res_1[offset + i];

auto e = Queue.submit([&](sycl::handler &cgh) {
auto a = buf_1.get_access<sycl::access::mode::read>(cgh, size / 2, offset);
auto b = buf_2.get_access<sycl::access::mode::write>(cgh, size / 2);
auto a = buf_1.get_access<sycl::access_mode::read>(cgh, size / 2, offset);
auto b = buf_2.get_access<sycl::access_mode::write>(cgh, size / 2);
cgh.copy(a, b);
});
e.wait();
Expand Down Expand Up @@ -141,10 +141,10 @@ void check_copy_host_to_device(sycl::queue &Queue) {
expected_res_4[i][j] = expected_res_3[i + offset][j + offset];

e = Queue.submit([&](sycl::handler &cgh) {
auto a = buf_3.get_access<sycl::access::mode::read>(
auto a = buf_3.get_access<sycl::access_mode::read>(
cgh, {size / 2, size / 2}, {offset, offset});
auto b =
buf_4.get_access<sycl::access::mode::write>(cgh, {size / 2, size / 2});
buf_4.get_access<sycl::access_mode::write>(cgh, {size / 2, size / 2});
cgh.copy(a, b);
});
e.wait();
Expand Down Expand Up @@ -190,9 +190,9 @@ void check_copy_host_to_device(sycl::queue &Queue) {
expected_res_5[i + offset][j + offset][k + offset];

e = Queue.submit([&](sycl::handler &cgh) {
auto a = buf_5.get_access<sycl::access::mode::read>(
auto a = buf_5.get_access<sycl::access_mode::read>(
cgh, {size / 2, size / 2, size / 2}, {offset, offset, offset});
auto b = buf_6.get_access<sycl::access::mode::write>(
auto b = buf_6.get_access<sycl::access_mode::write>(
cgh, {size / 2, size / 2, size / 2});
cgh.copy(a, b);
});
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/Basic/buffer/reinterpret.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ void execute_kernel(sycl::queue &cmd_queue, std::vector<OldType> &data,

cmd_queue.submit([&](sycl::handler &cgh) {
auto rb_acc =
reinterpret_subbuf.template get_access<sycl::access::mode::write>(cgh);
reinterpret_subbuf.template get_access<sycl::access_mode::write>(cgh);
cgh.parallel_for<KernelName>(
reinterpret_subbuf.get_range(),
[=](sycl::id<dim> index) { rb_acc[index] = val; });
Expand All @@ -61,7 +61,7 @@ int main() {
sycl::buffer<unsigned int, 1> buf_i(r1);
auto buf_char = buf_i.reinterpret<unsigned char>(r2);
cmd_queue.submit([&](sycl::handler &cgh) {
auto acc = buf_char.get_access<sycl::access::mode::read_write>(cgh);
auto acc = buf_char.get_access<sycl::access_mode::read_write>(cgh);
cgh.parallel_for<class chars>(r2,
[=](sycl::id<1> i) { acc[i] = UCHAR_MAX; });
});
Expand All @@ -81,7 +81,7 @@ int main() {
sycl::buffer<unsigned int, 1> buf_1d(r1d);
auto buf_2d = buf_1d.reinterpret<unsigned int>(r2d);
cmd_queue.submit([&](sycl::handler &cgh) {
auto acc2d = buf_2d.get_access<sycl::access::mode::read_write>(cgh);
auto acc2d = buf_2d.get_access<sycl::access_mode::read_write>(cgh);
cgh.parallel_for<class ones>(r2d, [=](sycl::item<2> itemID) {
size_t i = itemID.get_id(0);
size_t j = itemID.get_id(1);
Expand Down Expand Up @@ -194,7 +194,7 @@ int main() {

cmd_queue.submit([&](sycl::handler &cgh) {
auto rb_acc =
reinterpret_subbuf.get_access<sycl::access::mode::write>(cgh);
reinterpret_subbuf.get_access<sycl::access_mode::write>(cgh);
cgh.parallel_for<class foo_3>(
reinterpret_subbuf.get_range(),
[=](sycl::id<1> index) { rb_acc[index] = 13; });
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/Basic/char_builtins.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -46,9 +46,9 @@ template <typename T> int doCharTest(const T *A, const T *B, T *C) {
buffer<T> BBuf(B, BufferSize);
buffer<T> CBuf(C, BufferSize);
Q.submit([&](handler &CGH) {
auto AAcc = ABuf.template get_access<access::mode::read>(CGH);
auto BAcc = BBuf.template get_access<access::mode::read>(CGH);
auto CAcc = CBuf.template get_access<access::mode::write>(CGH);
auto AAcc = ABuf.template get_access<access_mode::read>(CGH);
auto BAcc = BBuf.template get_access<access_mode::read>(CGH);
auto CAcc = CBuf.template get_access<access_mode::write>(CGH);
CGH.single_task<>([=]() {
CAcc[0] = clz(AAcc[0]);
CAcc[1] = ctz(AAcc[1]);
Expand All @@ -69,7 +69,7 @@ template <typename T> int doCharTest(const T *A, const T *B, T *C) {
{
buffer<T> ABuf(A, BufferSize);
Q.submit([&](handler &CGH) {
auto AAcc = ABuf.template get_access<access::mode::read_write>(CGH);
auto AAcc = ABuf.template get_access<access_mode::read_write>(CGH);
local_accessor<T, 1> Local(range<1>{WorkGroupSize}, CGH);

nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)};
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/event.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -70,7 +70,7 @@ int main() {
float Scalar = 2.0;
sycl::buffer<float, 1> Buf(&Data, sycl::range<1>(1));
auto Event = Queue.submit([&](sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
auto Acc = Buf.get_access<sycl::access_mode::read_write>(cgh);

cgh.single_task<class test_event>([=]() { Acc[0] = Scalar; });
});
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/fill_accessor.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ void CheckFillZeroDimAccessor(queue &Q, T Init, T Expected) {
{
buffer<T, Dims> Buffer(Data.data(), Range);
Q.submit([&](handler &CGH) {
accessor<T, 0, sycl::access::mode::write> Accessor(Buffer, CGH);
accessor<T, 0, sycl::access_mode::write> Accessor(Buffer, CGH);
CGH.fill(Accessor, Expected);
}).wait_and_throw();
}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -30,12 +30,12 @@ int main() {
sycl::queue q;
sycl::nd_range<1> NDR(sycl::range<1>{n}, sycl::range<1>{2});
q.submit([&](sycl::handler &cgh) {
sycl::accessor<int, 1, sycl::access::mode::write,
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
acc(buf.get_access<sycl::access::mode::write>(cgh));
sycl::accessor<int, 1, sycl::access::mode::write,
acc(buf.get_access<sycl::access_mode::write>(cgh));
sycl::accessor<int, 1, sycl::access_mode::write,
sycl::access::target::device>
results_acc(results_buf.get_access<sycl::access::mode::write>(cgh));
results_acc(results_buf.get_access<sycl::access_mode::write>(cgh));
cgh.parallel_for<class NdItemTest>(NDR, [=](auto nd_i) {
static_assert(std::is_same<decltype(nd_i), sycl::nd_item<1>>::value,
"lambda arg type is unexpected");
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/Basic/group_async_copy_legacy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -114,8 +114,8 @@ template <typename T> int test(size_t Stride) {
// In must be read_write for legacy multi_ptr as it would otherwise produce
// a const multi_ptr which was not a valid source argument to
// async_work_group_copy in SYCL 1.2.1.
auto In = InBuf.template get_access<access::mode::read_write>(CGH);
auto Out = OutBuf.template get_access<access::mode::write>(CGH);
auto In = InBuf.template get_access<access_mode::read_write>(CGH);
auto Out = OutBuf.template get_access<access_mode::write>(CGH);
local_accessor<T, 1> Local(range<1>{WorkGroupSize}, CGH);

nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)};
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/Basic/handler/handler_copy_with_offset.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -19,8 +19,8 @@
#include <vector>

using namespace sycl;
constexpr access::mode read = access::mode::read;
constexpr access::mode write = access::mode::write;
constexpr access_mode read = access_mode::read;
constexpr access_mode write = access_mode::write;
constexpr access::target ondevice = access::target::device;

int main() {
Expand Down
23 changes: 11 additions & 12 deletions sycl/test-e2e/Basic/host-task-dependency.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -27,10 +27,9 @@ struct Context {

S::event HostTask_CopyBuf1ToBuf2(Context *Ctx) {
S::event Event = Ctx->Queue.submit([&](S::handler &CGH) {
S::host_accessor<int, 1, S::access::mode::read> CopierSrcAcc(Ctx->Buf1,
S::host_accessor<int, 1, S::access_mode::read> CopierSrcAcc(Ctx->Buf1, CGH);
S::host_accessor<int, 1, S::access_mode::write> CopierDstAcc(Ctx->Buf2,
CGH);
S::host_accessor<int, 1, S::access::mode::write> CopierDstAcc(Ctx->Buf2,
CGH);

auto CopierHostTask = [=] {
for (size_t Idx = 0; Idx < CopierDstAcc.size(); ++Idx)
Expand All @@ -54,29 +53,29 @@ S::event HostTask_CopyBuf1ToBuf2(Context *Ctx) {
void Thread1Fn(Context *Ctx) {
// 0. initialize resulting buffer with apriori wrong result
{
S::host_accessor<int, 1, S::access::mode::write> Acc(Ctx->Buf1);
S::host_accessor<int, 1, S::access_mode::write> Acc(Ctx->Buf1);

for (size_t Idx = 0; Idx < Acc.size(); ++Idx)
Acc[Idx] = -1;
}

{
S::host_accessor<int, 1, S::access::mode::write> Acc(Ctx->Buf2);
S::host_accessor<int, 1, S::access_mode::write> Acc(Ctx->Buf2);

for (size_t Idx = 0; Idx < Acc.size(); ++Idx)
Acc[Idx] = -2;
}

{
S::host_accessor<int, 1, S::access::mode::write> Acc(Ctx->Buf3);
S::host_accessor<int, 1, S::access_mode::write> Acc(Ctx->Buf3);

for (size_t Idx = 0; Idx < Acc.size(); ++Idx)
Acc[Idx] = -3;
}

// 1. submit task writing to buffer 1
Ctx->Queue.submit([&](S::handler &CGH) {
S::accessor<int, 1, S::access::mode::write, S::access::target::device>
S::accessor<int, 1, S::access_mode::write, S::access::target::device>
GeneratorAcc(Ctx->Buf1, CGH);

auto GeneratorKernel = [GeneratorAcc] {
Expand All @@ -92,9 +91,9 @@ void Thread1Fn(Context *Ctx) {

// 3. submit simple task to move data between two buffers
Ctx->Queue.submit([&](S::handler &CGH) {
S::accessor<int, 1, S::access::mode::read, S::access::target::device>
SrcAcc(Ctx->Buf2, CGH);
S::accessor<int, 1, S::access::mode::write, S::access::target::device>
S::accessor<int, 1, S::access_mode::read, S::access::target::device> SrcAcc(
Ctx->Buf2, CGH);
S::accessor<int, 1, S::access_mode::write, S::access::target::device>
DstAcc(Ctx->Buf3, CGH);

CGH.depends_on(HostTaskEvent);
Expand All @@ -109,7 +108,7 @@ void Thread1Fn(Context *Ctx) {

// 4. check data in buffer #3
{
S::host_accessor<int, 1, S::access::mode::read> Acc(Ctx->Buf3);
S::host_accessor<int, 1, S::access_mode::read> Acc(Ctx->Buf3);

bool Failure = false;

Expand Down Expand Up @@ -154,7 +153,7 @@ void test() {

// 3. check via host accessor that buf 2 contains valid data
{
S::host_accessor<int, 1, S::access::mode::read> ResultAcc(Ctx.Buf2);
S::host_accessor<int, 1, S::access_mode::read> ResultAcc(Ctx.Buf2);

bool Failure = false;
for (size_t Idx = 0; Idx < ResultAcc.size(); ++Idx) {
Expand Down
Loading
Loading