Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[SYCL][LIT]Update kernel-functor test #212

Merged
merged 2 commits into from
Apr 2, 2021
Merged

[SYCL][LIT]Update kernel-functor test #212

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
29f65e9
[SYCL][LIT]Update kernel-functor test
srividya-sundaram Mar 31, 2021
81aad4e
Update kernel_functor.cpp
srividya-sundaram Apr 1, 2021
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
36 changes: 5 additions & 31 deletions SYCL/Functor/kernel_functor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,8 @@
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out

//==--- kernel_functor.cpp - Functors as SYCL kernel test ------------------==//
//==--- kernel_functor.cpp -
// This test illustrates defining kernels as named function objects (functors)
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
Expand All @@ -20,27 +21,6 @@ constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;

// Case 1:
// - functor class is defined in an anonymous namespace
// - the '()' operator:
// * does not have parameters (to be used in 'single_task').
// * has the 'const' qualifier
namespace {
class Functor1 {
public:
Functor1(
int X_,
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &Acc_)
: X(X_), Acc(Acc_) {}

void operator()() const { Acc[0] += X; }

private:
int X;
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
};
} // namespace

// Case 2:
// - functor class is defined in a namespace
// - the '()' operator:
// * does not have parameters (to be used in 'single_task').
Expand All @@ -62,7 +42,7 @@ class Functor2 {
};
} // namespace ns

// Case 3:
// Case 2:
// - functor class is templated and defined in the translation unit scope
// - the '()' operator:
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
Expand All @@ -80,7 +60,7 @@ template <typename T> class TmplFunctor {
cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> Acc;
};

// Case 4:
// Case 3:
// - functor class is templated and defined in the translation unit scope
// - the '()' operator:
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
Expand All @@ -105,12 +85,6 @@ int foo(int X) {
cl::sycl::queue Q;
cl::sycl::buffer<int, 1> Buf(A, 1);

Q.submit([&](cl::sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
Functor1 F(X, Acc);

cgh.single_task(F);
});
Q.submit([&](cl::sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
ns::Functor2 F(X, Acc);
Expand Down Expand Up @@ -169,7 +143,7 @@ template <typename T> T bar(T X) {
int main() {
const int Res1 = foo(10);
const int Res2 = bar(10);
const int Gold1 = 40;
const int Gold1 = 30;
const int Gold2 = 80;

assert(Res1 == Gold1);
Expand Down