Skip to content

[SYCL] Make spec_constant default ctor public and available on host #2488

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Sep 17, 2020
Merged

[SYCL] Make spec_constant default ctor public and available on host #2488

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
9 changes: 3 additions & 6 deletions sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,17 +32,14 @@ class spec_const_error : public compile_program_error {
};

template <typename T, typename ID = T> class spec_constant {
private:
// Implementation defined constructor.
#ifdef __SYCL_DEVICE_ONLY__
public:
spec_constant() {}

private:
#else
spec_constant(T Cst) : Val(Cst) {}
#endif
#ifndef __SYCL_DEVICE_ONLY__
// Implementation defined constructor.
spec_constant(T Cst) : Val(Cst) {}

T Val;
#endif
friend class cl::sycl::program;
Expand Down
51 changes: 47 additions & 4 deletions sycl/test/spec_const/spec_const_hw.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,16 +23,17 @@

class MyInt32Const;
class MyFloatConst;
class MyConst;

using namespace sycl;

class KernelAAAi;
class KernelBBBf;

int val = 10;
int global_val = 10;

// Fetch a value at runtime.
int get_value() { return val; }
int get_value() { return global_val; }

float foo(
const cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst>
Expand All @@ -49,8 +50,22 @@ struct SCWrapper {
cl::sycl::ONEAPI::experimental::spec_constant<int, class sc_name2> SC2;
};

// MyKernel is used to test default constructor
using AccT = sycl::accessor<int, 1, sycl::access::mode::write>;
using ScT = sycl::ONEAPI::experimental::spec_constant<int, MyConst>;

struct MyKernel {
MyKernel(AccT &Acc) : Acc(Acc) {}

void setConst(ScT Sc) { this->Sc = Sc; }

void operator()() const { Acc[0] = Sc.get(); }
AccT Acc;
ScT Sc;
};

int main(int argc, char **argv) {
val = argc + 16;
global_val = argc + 16;

cl::sycl::queue q(default_selector{}, [](exception_list l) {
for (auto ep : l) {
Expand All @@ -68,10 +83,11 @@ int main(int argc, char **argv) {

std::cout << "Running on " << q.get_device().get_info<info::device::name>()
<< "\n";
std::cout << "val = " << val << "\n";
std::cout << "global_val = " << global_val << "\n";
cl::sycl::program program1(q.get_context());
cl::sycl::program program2(q.get_context());
cl::sycl::program program3(q.get_context());
cl::sycl::program program4(q.get_context());

int goldi = (int)get_value();
// TODO make this floating point once supported by the compiler
Expand All @@ -83,22 +99,30 @@ int main(int argc, char **argv) {
cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst> f32 =
program2.set_spec_constant<MyFloatConst>(goldf);

cl::sycl::ONEAPI::experimental::spec_constant<int, MyConst> sc =
program4.set_spec_constant<MyConst>(goldi);

program1.build_with_kernel_type<KernelAAAi>();
// Use an option (does not matter which exactly) to test different internal
// SYCL RT execution path
program2.build_with_kernel_type<KernelBBBf>("-cl-fast-relaxed-math");

SCWrapper W(program3);
program3.build_with_kernel_type<class KernelWrappedSC>();

program4.build_with_kernel_type<MyKernel>();

int goldw = 6;

std::vector<int> veci(1);
std::vector<float> vecf(1);
std::vector<int> vecw(1);
std::vector<int> vec(1);
try {
cl::sycl::buffer<int, 1> bufi(veci.data(), veci.size());
cl::sycl::buffer<float, 1> buff(vecf.data(), vecf.size());
cl::sycl::buffer<int, 1> bufw(vecw.data(), vecw.size());
cl::sycl::buffer<int, 1> buf(vec.data(), vec.size());

q.submit([&](cl::sycl::handler &cgh) {
auto acci = bufi.get_access<cl::sycl::access::mode::write>(cgh);
Expand All @@ -123,6 +147,19 @@ int main(int argc, char **argv) {
program3.get_kernel<KernelWrappedSC>(),
[=]() { accw[0] = W.SC1.get() + W.SC2.get(); });
});
// Check spec_constant default construction with subsequent initialization
q.submit([&](cl::sycl::handler &cgh) {
auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
// Specialization constants specification says:
// cl::sycl::experimental::spec_constant is default constructible,
// although the object is not considered initialized until the result of
// the call to cl::sycl::program::set_spec_constant is assigned to it.
MyKernel Kernel(acc); // default construct inside MyKernel instance
Kernel.setConst(sc); // initialize to sc, returned by set_spec_constant

cgh.single_task<MyKernel>(program4.get_kernel<MyKernel>(), Kernel);
});

} catch (cl::sycl::exception &e) {
std::cout << "*** Exception caught: " << e.what() << "\n";
return 1;
Expand All @@ -146,6 +183,12 @@ int main(int argc, char **argv) {
std::cout << "*** ERROR: " << valw << " != " << goldw << "(gold)\n";
passed = false;
}
int val = vec[0];

if (val != goldi) {
std::cout << "*** ERROR: " << val << " != " << goldi << "(gold)\n";
passed = false;
}
std::cout << (passed ? "passed\n" : "FAILED\n");
return passed ? 0 : 1;
}