From 9b01f2d12ac20af0664dbd1814255fe3546682df Mon Sep 17 00:00:00 2001 From: Noah Stuertz Date: Tue, 4 Feb 2025 08:12:42 +0100 Subject: [PATCH 1/6] Added Sycl statements for loops, buffer, kernel and accessors --- src/cpp.cc | 45 +++++++++++++++++++++++++++++ src/cpp.hh | 5 ++++ src/printer.cc | 6 ++++ src/printer.hh | 6 ++++ src/statement.cc | 12 ++++++++ src/statement.hh | 73 ++++++++++++++++++++++++++++++++++++++++++++++++ 6 files changed, 147 insertions(+) diff --git a/src/cpp.cc b/src/cpp.cc index ec62452d4..75558051c 100644 --- a/src/cpp.cc +++ b/src/cpp.cc @@ -105,6 +105,51 @@ void Printer::Cpp::print(const Statement::For &stmt) { stream << stmt.statements; } +void Printer::Cpp::print(const Statement::SYCL_Host_Accessor_Decl &stmt) { + assert(stmt.name); + + stream << indent() << "sycl::host_accessor " << stmt.name->name << "_hacc{" << stmt.name->name << "};" << endl; +} + + +void Printer::Cpp::print(const Statement::SYCL_Buffer_Decl &stmt) { + assert(stmt.type); + assert(stmt.dimension); + assert(stmt.name); + assert(stmt.value); + + stream << indent() << + "sycl::buffer<" << *stmt.type << "," << stmt.dimension << "> " << + *stmt.name << "(sycl::range<" << stmt.dimension << ">" << + "(" << *stmt.value->name << "));" << endl; +} + +void Printer::Cpp::print(const Statement::SYCL_Accessor_Decl &stmt) { + assert(stmt.variable); + assert(stmt.conext); + + stream << indent() << "auto " << *stmt.variable->name + "_acc = sycl::accessor{" + << *stmt.variable->name << ", " << *stmt.context->name << ", "; + if (*stmt.write && *stmt.read) { + stream << "sycl::access_mode::read_write"; + } else if (*stmt.write) { + stream << "sycl::access_mode::write"; + } else { + stream << "sycl::access_mode::read"; + } + stream << "};" << endl; +} + +void Printer::Cpp::print(const Statement::SYCL_Submit_Kernel &stmt) { + assert(stmt.queue); + assert(stmt.context); + + stream << indent() << + *stmt.queue->name << ".submit([&]sycl::handler &" << *stmt.context->name << ") "; + stream << stmt.statements; + stream << ");" << endl; +} + void Printer::Cpp::print(const Statement::While &stmt) { stream << indent() << "while(" << stmt.expr() << ")\n"; stream << stmt.statements; diff --git a/src/cpp.hh b/src/cpp.hh index 60dbe3f83..20a6e903f 100644 --- a/src/cpp.hh +++ b/src/cpp.hh @@ -150,6 +150,11 @@ class Cpp : public Base { void print(const Statement::Hash_Decl &stmt); void print(const Statement::Marker_Decl &stmt); + void print(const Statement::SYCL_Accessor_Decl &stmt); + void print(const Statement::SYCL_Submit_Kernel &stmt); + void print(const Statement::SYCL_Buffer_Decl &stmt); + void print(const Statement::SYCL_Host_Accessor_Decl &stmt); + void print(const Fn_Def &fn_def); void print(const Operator &op); diff --git a/src/printer.cc b/src/printer.cc index d4777963a..e67d7d96d 100644 --- a/src/printer.cc +++ b/src/printer.cc @@ -78,6 +78,12 @@ void Printer::Base::print(const Statement::Table_Decl &stmt) {} void Printer::Base::print(const std::list &stmts) {} +void Printer::Base::print(const Statement::SYCL_Buffer_Decl &stmt) {} +void Printer::Base::print(const Statement::SYCL_Submit_Kernel &stmt) {} +void Printer::Base::print(const Statement::SYCL_Accessor_Decl &stmt) {} +void Printer::Base::print(const Statement::SYCL_Host_Accessor_Decl &stmt) {} + + void Printer::Base::print(const Type::List &t) {} void Printer::Base::print(const Type::Tuple &t) {} void Printer::Base::print(const Type::TupleDef &t) {} diff --git a/src/printer.hh b/src/printer.hh index ef3436f4b..ed25fd31a 100644 --- a/src/printer.hh +++ b/src/printer.hh @@ -124,6 +124,12 @@ class Base { virtual void print(const Statement::Table_Decl &stmt); + virtual void print(const Statement::SYCL_Buffer_Decl &stmt); + virtual void print(const Statement::SYCL_Submit_Kernel &stmt); + virtual void print(const Statement::SYCL_Accessor_Decl &stmt); + virtual void print(const Statement::SYCL_Host_Accessor_Decl &stmt); + + virtual void print(const Expr::Base &); virtual void print(const Type::Base &); virtual void print(const Var_Acc::Base &); diff --git a/src/statement.cc b/src/statement.cc index e911c3041..05331fbbd 100644 --- a/src/statement.cc +++ b/src/statement.cc @@ -71,6 +71,18 @@ Statement::Var_Decl::Var_Decl(::Type::Base *t, std::string *n, Expr::Base *e) : Base(VAR_DECL), type(t), name(n), rhs(e) { } +Statement::SYCL_Buffer_Decl::SYCL_Buffer_Decl(::Type::Base *t, int d, std::string *n, Var_Decl *v) + : Base(BUFFER_DECL), type(t), dimension(d), name(n), value(v) { +} + +Statement::SYCL_Accessor_Decl::SYCL_Accessor_Decl(Var_Decl *v, Var_Decl *c, bool *r, bool *w) + : Base(VAR_DECL), variable(v), context(c), read(r), write(w) { +} + +Statement::SYCL_Host_Accessor_Decl::SYCL_Host_Accessor_Decl(Var_Decl *n) + : Base(VAR_DECL), name(n) { + + } Statement::Var_Decl *Statement::Var_Decl::clone() const { Var_Decl *ret = new Var_Decl(*this); diff --git a/src/statement.hh b/src/statement.hh index e20a432da..3b2ce20a1 100644 --- a/src/statement.hh +++ b/src/statement.hh @@ -194,6 +194,79 @@ class Increase : public Base { }; +/** + * @brief Declare a host accessor for Variables + * @example sycl::host_accessor result{results}; + * @param name the name of the variable you want a host accessor for + */ +class SYCL_Host_Accessor_Decl : public Base { + public: + Var_Decl *name; + + void print(Printer::Base &p) const; + + SYCL_Host_Accessor_Decl(Var_Decl *n); +}; +/** + * @example auto aResult = sycl::accessor{results, cgh, sycl::read_write}; + * @brief Declare a Accessor for SYCL Kernels + * @param variable Var_Decl name and variable to create accessor to + * @param context Context defaults to cgh + * @param access_mode Access Mode (Read or Write) + */ +class SYCL_Accessor_Decl : public Base { + + public: + Var_Decl *variable; + Var_Decl *context; + bool *read; + bool *write; + + void print(Printer::Base &p) const; + + SYCL_Accessor_Decl(Var_Decl *v, Var_Decl *c, bool *r, bool *w); +}; + +/** + * @example q.submit([&](sycl::handler &cgh) { ... }); + * @brief Submit Kernel + * @param queue The Queue to add the Kernel to + * @param context The Context Handler for the Kernel + */ +class SYCL_Submit_Kernel : public Block_Base { + public: + Var_Decl *queue; + Var_Decl *context; + + void print(Printer::Base &p) const; + + SYCL_Submit_Kernel(Var_Decl *q, Var_Decl *c); +}; + +/** + * @example sycl::buffer results(sycl::range<1>(n*m)); + * @brief A function to summarize and create an buffer + * + * @param type What type the buffer should hold + * @param dimension + * @param name + * @param size + */ + class SYCL_Buffer_Decl : public Base { + public: + ::Type::Base *type; + int dimension; + std::string *name; + Var_Decl *value; + + SYCL_Buffer_Decl(::Type::Base *t, int d, std::string *n, Var_Decl *v); + + SYCL_Buffer_Decl(::Type::Base *t, int d, Var_Decl *v, const std::string &n) + : Base(BUFFER_DECL), type(t), dimension(d), value(v) { name = new std::string(n); } + + void print(Printer::Base &p) const; + }; + class Var_Decl : public Base { private: Bool use_as_itr; From 7463937c093ce9db90b276ec81b5cc5d019e3e83 Mon Sep 17 00:00:00 2001 From: Noah Stuertz Date: Wed, 29 Jan 2025 20:04:21 +0100 Subject: [PATCH 2/6] Using Example Code in cyk.cc --- src/cyk.cc | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/src/cyk.cc b/src/cyk.cc index fc61bffaa..a2a94245e 100644 --- a/src/cyk.cc +++ b/src/cyk.cc @@ -1299,7 +1299,21 @@ Fn_Def *print_CYK(const AST &ast) { } } } - fn_cyk->stmts.push_back(new Statement::CustomCode("#pragma omp parallel")); + + int dimension = 1; + std::string name = "test"; + std::string value = "test_value"; + + fn_cyk->stmts.push_back(new Statement::SYCL_Buffer_Decl(new Type::Int, dimension, name, value)); + Statement::Var_Decl *queue = new Statement::Var_Decl(new Type::External("sycl::queue"), "q"); + + fn_cyk->stmts.push_back(queue); + + Statement::SYCL_Submit_Kernel *blk_sycl = new Statement::SYCL_Submit_Kernel(queue, new Statement::Var_Decl(new Type::External("sycl::handler&"), "cgh")); + + bool* test = new bool(true); + blk_sycl->statements.push_back(new Statement::SYCL_Accessor_Decl(value, value, test, test)); + Statement::Block *blk_parallel = new Statement::Block(); if (ast.checkpoint && ast.checkpoint->cyk) { From 0245d2dbd6d27f2cab69e94ed149bfef2434aa16 Mon Sep 17 00:00:00 2001 From: Noah Stuertz Date: Wed, 5 Feb 2025 17:08:35 +0100 Subject: [PATCH 3/6] fixed host_accessor statement --- src/statement.cc | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/src/statement.cc b/src/statement.cc index 05331fbbd..a721537ce 100644 --- a/src/statement.cc +++ b/src/statement.cc @@ -94,6 +94,10 @@ Statement::Var_Decl *Statement::Var_Decl::clone() const { } +void Statement::SYCL_Host_Accessor_Decl::print(Printer::Base &p) const { + p.print(*this); +} + void Statement::Var_Decl::print(Printer::Base &p) const { p.print(*this); } From b41d62cfae72f85b881e297aaf17942c23b22c86 Mon Sep 17 00:00:00 2001 From: Noah Stuertz Date: Thu, 20 Mar 2025 12:52:21 +0100 Subject: [PATCH 4/6] fixed local linting errors --- src/cpp.cc | 24 ++++++++++++++++-------- src/cyk.cc | 14 ++++++++++---- src/statement.cc | 9 +++++---- src/statement.hh | 29 +++++++++++++++-------------- 4 files changed, 46 insertions(+), 30 deletions(-) diff --git a/src/cpp.cc b/src/cpp.cc index 75558051c..0470f39d9 100644 --- a/src/cpp.cc +++ b/src/cpp.cc @@ -108,7 +108,10 @@ void Printer::Cpp::print(const Statement::For &stmt) { void Printer::Cpp::print(const Statement::SYCL_Host_Accessor_Decl &stmt) { assert(stmt.name); - stream << indent() << "sycl::host_accessor " << stmt.name->name << "_hacc{" << stmt.name->name << "};" << endl; + stream << indent() + << "sycl::host_accessor " << stmt.name->name << "_hacc{" + << stmt.name->name << "};" + << endl; } @@ -118,18 +121,20 @@ void Printer::Cpp::print(const Statement::SYCL_Buffer_Decl &stmt) { assert(stmt.name); assert(stmt.value); - stream << indent() << + stream << indent() << "sycl::buffer<" << *stmt.type << "," << stmt.dimension << "> " << - *stmt.name << "(sycl::range<" << stmt.dimension << ">" << + *stmt.name << "(sycl::range<" << stmt.dimension << ">" << "(" << *stmt.value->name << "));" << endl; } void Printer::Cpp::print(const Statement::SYCL_Accessor_Decl &stmt) { assert(stmt.variable); assert(stmt.conext); - - stream << indent() << "auto " << *stmt.variable->name + "_acc = sycl::accessor{" + + stream << indent() << + "auto " << *stmt.variable->name + "_acc = sycl::accessor{" << *stmt.variable->name << ", " << *stmt.context->name << ", "; + if (*stmt.write && *stmt.read) { stream << "sycl::access_mode::read_write"; } else if (*stmt.write) { @@ -137,7 +142,8 @@ void Printer::Cpp::print(const Statement::SYCL_Accessor_Decl &stmt) { } else { stream << "sycl::access_mode::read"; } - stream << "};" << endl; + + stream << "};" << endl; } void Printer::Cpp::print(const Statement::SYCL_Submit_Kernel &stmt) { @@ -145,10 +151,12 @@ void Printer::Cpp::print(const Statement::SYCL_Submit_Kernel &stmt) { assert(stmt.context); stream << indent() << - *stmt.queue->name << ".submit([&]sycl::handler &" << *stmt.context->name << ") "; + *stmt.queue->name << + ".submit([&]sycl::handler &" << *stmt.context->name << ") "; + stream << stmt.statements; stream << ");" << endl; -} +} void Printer::Cpp::print(const Statement::While &stmt) { stream << indent() << "while(" << stmt.expr() << ")\n"; diff --git a/src/cyk.cc b/src/cyk.cc index a2a94245e..13117c08e 100644 --- a/src/cyk.cc +++ b/src/cyk.cc @@ -1304,15 +1304,21 @@ Fn_Def *print_CYK(const AST &ast) { std::string name = "test"; std::string value = "test_value"; - fn_cyk->stmts.push_back(new Statement::SYCL_Buffer_Decl(new Type::Int, dimension, name, value)); - Statement::Var_Decl *queue = new Statement::Var_Decl(new Type::External("sycl::queue"), "q"); + fn_cyk->stmts.push_back( + new Statement::SYCL_Buffer_Decl(new Type::Int, dimension, name, value)); + + Statement::Var_Decl *queue = new Statement::Var_Decl( + new Type::External("sycl::queue"), "q"); fn_cyk->stmts.push_back(queue); - Statement::SYCL_Submit_Kernel *blk_sycl = new Statement::SYCL_Submit_Kernel(queue, new Statement::Var_Decl(new Type::External("sycl::handler&"), "cgh")); + Statement::SYCL_Submit_Kernel *blk_sycl = new Statement::SYCL_Submit_Kernel( + queue, new Statement::Var_Decl( + new Type::External("sycl::handler&"), "cgh")); bool* test = new bool(true); - blk_sycl->statements.push_back(new Statement::SYCL_Accessor_Decl(value, value, test, test)); + blk_sycl->statements.push_back( + new Statement::SYCL_Accessor_Decl(value, value, test, test)); Statement::Block *blk_parallel = new Statement::Block(); diff --git a/src/statement.cc b/src/statement.cc index a721537ce..794ecc112 100644 --- a/src/statement.cc +++ b/src/statement.cc @@ -71,18 +71,19 @@ Statement::Var_Decl::Var_Decl(::Type::Base *t, std::string *n, Expr::Base *e) : Base(VAR_DECL), type(t), name(n), rhs(e) { } -Statement::SYCL_Buffer_Decl::SYCL_Buffer_Decl(::Type::Base *t, int d, std::string *n, Var_Decl *v) +Statement::SYCL_Buffer_Decl::SYCL_Buffer_Decl(::Type::Base *t, int d, + std::string *n, Var_Decl *v) : Base(BUFFER_DECL), type(t), dimension(d), name(n), value(v) { } -Statement::SYCL_Accessor_Decl::SYCL_Accessor_Decl(Var_Decl *v, Var_Decl *c, bool *r, bool *w) +Statement::SYCL_Accessor_Decl::SYCL_Accessor_Decl(Var_Decl *v, Var_Decl *c, + bool *r, bool *w) : Base(VAR_DECL), variable(v), context(c), read(r), write(w) { } Statement::SYCL_Host_Accessor_Decl::SYCL_Host_Accessor_Decl(Var_Decl *n) : Base(VAR_DECL), name(n) { - - } +} Statement::Var_Decl *Statement::Var_Decl::clone() const { Var_Decl *ret = new Var_Decl(*this); diff --git a/src/statement.hh b/src/statement.hh index 3b2ce20a1..7cc17cc1b 100644 --- a/src/statement.hh +++ b/src/statement.hh @@ -200,12 +200,12 @@ class Increase : public Base { * @param name the name of the variable you want a host accessor for */ class SYCL_Host_Accessor_Decl : public Base { - public: - Var_Decl *name; - + public: + Var_Decl *name; + void print(Printer::Base &p) const; - SYCL_Host_Accessor_Decl(Var_Decl *n); + explicit SYCL_Host_Accessor_Decl(Var_Decl *n); }; /** * @example auto aResult = sycl::accessor{results, cgh, sycl::read_write}; @@ -215,8 +215,7 @@ class SYCL_Host_Accessor_Decl : public Base { * @param access_mode Access Mode (Read or Write) */ class SYCL_Accessor_Decl : public Base { - - public: + public: Var_Decl *variable; Var_Decl *context; bool *read; @@ -234,7 +233,7 @@ class SYCL_Accessor_Decl : public Base { * @param context The Context Handler for the Kernel */ class SYCL_Submit_Kernel : public Block_Base { - public: + public: Var_Decl *queue; Var_Decl *context; @@ -252,20 +251,22 @@ class SYCL_Submit_Kernel : public Block_Base { * @param name * @param size */ - class SYCL_Buffer_Decl : public Base { - public: +class SYCL_Buffer_Decl : public Base { + public: ::Type::Base *type; int dimension; std::string *name; Var_Decl *value; - SYCL_Buffer_Decl(::Type::Base *t, int d, std::string *n, Var_Decl *v); + SYCL_Buffer_Decl(::Type::Base *t, int d, std::string *n, Var_Decl *v); - SYCL_Buffer_Decl(::Type::Base *t, int d, Var_Decl *v, const std::string &n) - : Base(BUFFER_DECL), type(t), dimension(d), value(v) { name = new std::string(n); } + SYCL_Buffer_Decl(::Type::Base *t, int d, Var_Decl *v, const std::string &n) + : Base(BUFFER_DECL), type(t), dimension(d), value(v) { + name = new std::string(n); + } - void print(Printer::Base &p) const; - }; + void print(Printer::Base &p) const; +}; class Var_Decl : public Base { private: From 34d8f6f4b4fb741abeaaf70be9638d90a5973900 Mon Sep 17 00:00:00 2001 From: Noah Stuertz Date: Thu, 20 Mar 2025 15:18:14 +0100 Subject: [PATCH 5/6] Fixed random missing lines and wrong types --- src/cyk.cc | 5 +++-- src/statement.cc | 30 +++++++++++++++++++++++------- src/statement.hh | 9 ++------- 3 files changed, 28 insertions(+), 16 deletions(-) diff --git a/src/cyk.cc b/src/cyk.cc index 13117c08e..0e433eb12 100644 --- a/src/cyk.cc +++ b/src/cyk.cc @@ -26,6 +26,7 @@ #include #include #include +#include "statement.hh" static const char *MUTEX = "mutex"; static const char *VARNAME_OuterLoop1 = "outer_loop_1_idx"; @@ -1302,10 +1303,10 @@ Fn_Def *print_CYK(const AST &ast) { int dimension = 1; std::string name = "test"; - std::string value = "test_value"; + Statement::Var_Decl *value = new Statement::Var_Decl(new Type::String, "test_value"); fn_cyk->stmts.push_back( - new Statement::SYCL_Buffer_Decl(new Type::Int, dimension, name, value)); + new Statement::SYCL_Buffer_Decl(new Type::Int, dimension, value, value)); Statement::Var_Decl *queue = new Statement::Var_Decl( new Type::External("sycl::queue"), "q"); diff --git a/src/statement.cc b/src/statement.cc index 794ecc112..070eff9ed 100644 --- a/src/statement.cc +++ b/src/statement.cc @@ -31,6 +31,7 @@ #include "statement.hh" +#include "statement/base.hh" #include "var_acc.hh" #include "expr.hh" @@ -71,19 +72,22 @@ Statement::Var_Decl::Var_Decl(::Type::Base *t, std::string *n, Expr::Base *e) : Base(VAR_DECL), type(t), name(n), rhs(e) { } -Statement::SYCL_Buffer_Decl::SYCL_Buffer_Decl(::Type::Base *t, int d, - std::string *n, Var_Decl *v) - : Base(BUFFER_DECL), type(t), dimension(d), name(n), value(v) { +Statement::SYCL_Buffer_Decl::SYCL_Buffer_Decl(::Type::Base *t, int d, Var_Decl *v, Var_Decl *n) + : Base(VAR_DECL), type(t), dimension(d), value(v), name(n) { } -Statement::SYCL_Accessor_Decl::SYCL_Accessor_Decl(Var_Decl *v, Var_Decl *c, - bool *r, bool *w) +Statement::SYCL_Accessor_Decl::SYCL_Accessor_Decl(Var_Decl *v, Var_Decl *c, bool *r, bool *w) : Base(VAR_DECL), variable(v), context(c), read(r), write(w) { } Statement::SYCL_Host_Accessor_Decl::SYCL_Host_Accessor_Decl(Var_Decl *n) - : Base(VAR_DECL), name(n) { -} + : Base(VAR_DECL), name(n) { + } + +Statement::SYCL_Submit_Kernel::SYCL_Submit_Kernel(Var_Decl *q, Var_Decl *c) + : Block_Base(BLOCK), queue(q), context(c) { + + } Statement::Var_Decl *Statement::Var_Decl::clone() const { Var_Decl *ret = new Var_Decl(*this); @@ -95,6 +99,18 @@ Statement::Var_Decl *Statement::Var_Decl::clone() const { } +void Statement::SYCL_Accessor_Decl::print(Printer::Base &p) const { + p.print(*this); +} + +void Statement::SYCL_Buffer_Decl::print(Printer::Base &p) const { +p.print(*this); +} + +void Statement::SYCL_Submit_Kernel::print(Printer::Base &p) const { + p.print(*this); +} + void Statement::SYCL_Host_Accessor_Decl::print(Printer::Base &p) const { p.print(*this); } diff --git a/src/statement.hh b/src/statement.hh index 7cc17cc1b..b468921b6 100644 --- a/src/statement.hh +++ b/src/statement.hh @@ -255,15 +255,10 @@ class SYCL_Buffer_Decl : public Base { public: ::Type::Base *type; int dimension; - std::string *name; + Var_Decl *name; Var_Decl *value; - SYCL_Buffer_Decl(::Type::Base *t, int d, std::string *n, Var_Decl *v); - - SYCL_Buffer_Decl(::Type::Base *t, int d, Var_Decl *v, const std::string &n) - : Base(BUFFER_DECL), type(t), dimension(d), value(v) { - name = new std::string(n); - } + SYCL_Buffer_Decl(::Type::Base *t, int d, Var_Decl *v, Var_Decl *n); void print(Printer::Base &p) const; }; From 15b0d260087d6028295300233b546b76ee2e411c Mon Sep 17 00:00:00 2001 From: Noah Stuertz Date: Mon, 24 Mar 2025 10:44:15 +0100 Subject: [PATCH 6/6] fixed final linting errors --- src/cyk.cc | 3 ++- src/statement.cc | 9 +++++---- 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/src/cyk.cc b/src/cyk.cc index 0e433eb12..dce3f5db9 100644 --- a/src/cyk.cc +++ b/src/cyk.cc @@ -1303,7 +1303,8 @@ Fn_Def *print_CYK(const AST &ast) { int dimension = 1; std::string name = "test"; - Statement::Var_Decl *value = new Statement::Var_Decl(new Type::String, "test_value"); + Statement::Var_Decl *value = new Statement::Var_Decl( + new Type::String, "test_value"); fn_cyk->stmts.push_back( new Statement::SYCL_Buffer_Decl(new Type::Int, dimension, value, value)); diff --git a/src/statement.cc b/src/statement.cc index 070eff9ed..c53efabf5 100644 --- a/src/statement.cc +++ b/src/statement.cc @@ -72,21 +72,22 @@ Statement::Var_Decl::Var_Decl(::Type::Base *t, std::string *n, Expr::Base *e) : Base(VAR_DECL), type(t), name(n), rhs(e) { } -Statement::SYCL_Buffer_Decl::SYCL_Buffer_Decl(::Type::Base *t, int d, Var_Decl *v, Var_Decl *n) +Statement::SYCL_Buffer_Decl::SYCL_Buffer_Decl( + ::Type::Base *t, int d, Var_Decl *v, Var_Decl *n) : Base(VAR_DECL), type(t), dimension(d), value(v), name(n) { } -Statement::SYCL_Accessor_Decl::SYCL_Accessor_Decl(Var_Decl *v, Var_Decl *c, bool *r, bool *w) +Statement::SYCL_Accessor_Decl::SYCL_Accessor_Decl( + Var_Decl *v, Var_Decl *c, bool *r, bool *w) : Base(VAR_DECL), variable(v), context(c), read(r), write(w) { } Statement::SYCL_Host_Accessor_Decl::SYCL_Host_Accessor_Decl(Var_Decl *n) - : Base(VAR_DECL), name(n) { + : Base(VAR_DECL), name(n) { } Statement::SYCL_Submit_Kernel::SYCL_Submit_Kernel(Var_Decl *q, Var_Decl *c) : Block_Base(BLOCK), queue(q), context(c) { - } Statement::Var_Decl *Statement::Var_Decl::clone() const {