From 2a26824aa298845bd49d50148c12e04cbc369414 Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Tue, 10 Jan 2023 12:28:45 +0530 Subject: [PATCH 1/8] Add sycl backend --- cgen/__init__.py | 10 +++ cgen/sycl.py | 219 +++++++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 229 insertions(+) create mode 100644 cgen/sycl.py diff --git a/cgen/__init__.py b/cgen/__init__.py index 9bd7522..98477fb 100644 --- a/cgen/__init__.py +++ b/cgen/__init__.py @@ -413,7 +413,17 @@ def struct_format(self): mapper_method = "map_function_declaration" # }}} +class Lamda(NestedDeclarator): + def __init__(self, capture_clause,subdecl ,arg_decls): + self.capture_clause = capture_clause + self.arg_decls = arg_decls + NestedDeclarator.__init__(self, subdecl) + + def get_decl_pair(self): + sub_tp, sub_decl = self.subdecl.get_decl_pair() + return sub_tp, f"{sub_decl} = [{self.capture_clause}]({'{}'.format(', '.join(ad.inline() for ad in self.arg_decls))})" + mapper_method = "map_lamda" # {{{ struct-like diff --git a/cgen/sycl.py b/cgen/sycl.py new file mode 100644 index 0000000..e35e5ff --- /dev/null +++ b/cgen/sycl.py @@ -0,0 +1,219 @@ +__copyright__ = "Copyright (C) 2011-20 Andreas Kloeckner" + +__license__ = """ +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +""" + +import numpy as np + +from cgen import \ + Declarator, \ + DeclSpecifier, \ + NestedDeclarator, \ + Value,\ + Pointer,\ + Extern,\ +NamespaceQualifier,\ + Generable + + +def dtype_to_cltype(dtype): + if dtype is None: + raise ValueError("dtype may not be None") + + dtype = np.dtype(dtype) + if dtype == np.int64: + return "long" + elif dtype == np.uint64: + return "ulong" + elif dtype == np.int32: + return "int" + elif dtype == np.uint32: + return "uint" + elif dtype == np.int16: + return "short" + elif dtype == np.uint16: + return "ushort" + elif dtype == np.int8: + return "char" + elif dtype == np.uint8: + return "uchar" + elif dtype == np.float32: + return "float" + elif dtype == np.float64: + return "double" + else: + raise ValueError(f"unable to map dtype '{dtype}'") + + +# {{{ kernel + +class SYCLKernel(DeclSpecifier): + def __init__(self, subdecl,ndim): + self.ndim=ndim + subdecl.arg_decls.append(Value("sycl::queue","queue_")) + subdecl.arg_decls.append(Value("sycl::nd_range<{}> ".format(self.ndim),"range_")) + DeclSpecifier.__init__(self, subdecl,"") + + + mapper_method = "map_cl_kernel" + +# }}} + + +# {{{ kernel args + +class CLConstant(DeclSpecifier): + def __init__(self, subdecl): + DeclSpecifier.__init__(self, subdecl, "__constant") + + mapper_method = "map_cl_constant" + + +class CLLocal(DeclSpecifier): + def __init__(self, subdecl): + DeclSpecifier.__init__(self, subdecl, "__local") + + mapper_method = "map_cl_local" + + +class CLGlobal(DeclSpecifier): + def __init__(self, subdecl): + DeclSpecifier.__init__(self, subdecl, "__global") + + mapper_method = "map_cl_global" + + +class CLImage(Value): + def __init__(self, dims, mode, name): + if mode == "r": + spec = "__read_only" + elif mode == "w": + spec = "__write_only" + else: + raise ValueError("mode must be one of 'r' or 'w'") + + Value.__init__(self, f"{spec} image{dims}d_t", name) + + mapper_method = "map_cl_image" + +# }}} + + +# {{{ function attributes + +class CLVecTypeHint(NestedDeclarator): + def __init__(self, subdecl, dtype=None, count=None, type_str=None): + if (dtype is None) != (count is None): + raise ValueError("dtype and count must always be " + "specified together") + + if (dtype is None and type_str is None) or \ + (dtype is not None and type_str is not None): + raise ValueError("exactly one of dtype and type_str must be specified") + + if type_str is None: + self.type_str = dtype_to_cltype(dtype)+str(count) + else: + self.type_str = type_str + + super().__init__(subdecl) + + def get_decl_pair(self): + sub_tp, sub_decl = self.subdecl.get_decl_pair() + return sub_tp, ("__attribute__ ((vec_type_hint({}))) {}".format( + sub_decl, self.type_str)) + + mapper_method = "map_cl_vec_type_hint" + + +class _SYCLWorkGroupSizeDeclarator(Generable): + def __init__(self, dim): + + while len(dim) < 3: + dim = dim + (1,) + self.dim = dim + self.decl='extern "C" [[sycl::reqd_work_group_size({})]]'.format(", ".join(str(d) for d in self.dim)) + + def generate(self): + yield self.decl + +class CLWorkGroupSizeHint(_SYCLWorkGroupSizeDeclarator): + """ + See Sec 6.7.2 of OpenCL 2.0 spec, Version V2.2-11. + """ + + mapper_method = "map_cl_workgroup_size_hint" + + +class SYCLRequiredWorkGroupSize(_SYCLWorkGroupSizeDeclarator): + """ + See Sec 6.7.2 of OpenCL 2.0 spec, Version V2.2-11. + """ + mapper_method = "map_cl_required_wokgroup_size" + +# }}} + + +# {{{ vector PODs + +class CLVectorPOD(Declarator): + def __init__(self, dtype, count, name): + self.dtype = np.dtype(dtype) + self.count = count + self.name = name + + def get_decl_pair(self): + return [dtype_to_cltype(self.dtype)+str(self.count)], self.name + + def struct_maker_code(self, name): + return name + + def struct_format(self): + return str(self.count)+self.dtype.char + + def alignment_requirement(self): + from struct import calcsize + return calcsize(self.struct_format()) + + def default_value(self): + return [0]*self.count + + mapper_method = "map_cl_vector_pod" + +# }}} + +# vim: fdm=marker + +class SYCLBody(Generable): + def __init__(self, body,ndim): + """Initialize a function definition. *fdecl* is expected to be + a :class:`FunctionDeclaration` instance, while *body* is a + :class:`Block`. + """ + + self.ndim=ndim + self.upper = "queue_.submit([&](sycl::handler &h) {"+"\n h.parallel_for(range_, [=](sycl::nd_item<{}> item)".format(self.ndim) + self.body = body + self.lower=");\n }).wait();" + + def generate(self): + yield self.upper + yield from self.body.generate() + yield self.lower + + mapper_method = "map_function_body" From fc14c9e1a512b90d6fabf1955540882d98e2ec19 Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Tue, 7 Mar 2023 16:34:15 +0530 Subject: [PATCH 2/8] Add SYCL accessor --- cgen/sycl.py | 31 ++++++++++++++++++++++++++++--- 1 file changed, 28 insertions(+), 3 deletions(-) diff --git a/cgen/sycl.py b/cgen/sycl.py index e35e5ff..38c5461 100644 --- a/cgen/sycl.py +++ b/cgen/sycl.py @@ -84,11 +84,17 @@ def __init__(self, subdecl): mapper_method = "map_cl_constant" -class CLLocal(DeclSpecifier): +class SYCLLocal(NestedDeclarator): def __init__(self, subdecl): - DeclSpecifier.__init__(self, subdecl, "__local") + print(subdecl) + self.type=subdecl.get_type() + self.subdecl=subdecl - mapper_method = "map_cl_local" + def get_decl_pair(self): + sub_tp,sub_decl=self.subdecl.get_decl_pair() + return [f"sycl::local_accessor<{self.type}>"],sub_decl + + mapper_method = "map_sycl_local" class CLGlobal(DeclSpecifier): @@ -198,7 +204,26 @@ def default_value(self): # }}} # vim: fdm=marker +class SYCLAccessor(NestedDeclarator): + def __init__(self, subdecl,handler,count=None): + NestedDeclarator.__init__(self, subdecl) + self.handler=handler + self.count=count + sub_tp, sub_decl = subdecl.get_decl_pair() + self.type=sub_tp[0] + self.sub_decl=sub_decl + + def get_decl_pair(self): + if self.count is None: + return [f"sycl::accessor<{ self.type}>"], f"{self.sub_decl}({self.handler})" + else: + return [f"sycl::accessor<{ self.type}>"], f"{self.sub_decl}(sycl::range<1>({self.count}),{self.handler})" + + def get_type(self): + return self.type + mapper_method = "map_sycl_accessor" + class SYCLBody(Generable): def __init__(self, body,ndim): """Initialize a function definition. *fdecl* is expected to be From a63d918607b3f5e14595a6d704ef5adcebbc6424 Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Wed, 8 Mar 2023 17:22:50 +0530 Subject: [PATCH 3/8] Add SYCLparallel_for, SYCLQueueSubmit functions --- cgen/__init__.py | 8 +++----- cgen/sycl.py | 39 +++++++++++++++++++++++++-------------- 2 files changed, 28 insertions(+), 19 deletions(-) diff --git a/cgen/__init__.py b/cgen/__init__.py index 98477fb..8b89f90 100644 --- a/cgen/__init__.py +++ b/cgen/__init__.py @@ -413,15 +413,13 @@ def struct_format(self): mapper_method = "map_function_declaration" # }}} -class Lamda(NestedDeclarator): - def __init__(self, capture_clause,subdecl ,arg_decls): +class Lamda(Declarator): + def __init__(self, capture_clause ,arg_decls): self.capture_clause = capture_clause self.arg_decls = arg_decls - NestedDeclarator.__init__(self, subdecl) def get_decl_pair(self): - sub_tp, sub_decl = self.subdecl.get_decl_pair() - return sub_tp, f"{sub_decl} = [{self.capture_clause}]({'{}'.format(', '.join(ad.inline() for ad in self.arg_decls))})" + return [f"[{self.capture_clause}]"], f"({'{}'.format(', '.join(ad.inline() for ad in self.arg_decls))})" mapper_method = "map_lamda" diff --git a/cgen/sycl.py b/cgen/sycl.py index 38c5461..7f66c0d 100644 --- a/cgen/sycl.py +++ b/cgen/sycl.py @@ -27,8 +27,11 @@ Value,\ Pointer,\ Extern,\ -NamespaceQualifier,\ - Generable + NamespaceQualifier,\ + Generable,\ + FunctionBody,\ + Lamda,\ + Block def dtype_to_cltype(dtype): @@ -225,20 +228,28 @@ def get_type(self): class SYCLBody(Generable): - def __init__(self, body,ndim): - """Initialize a function definition. *fdecl* is expected to be - a :class:`FunctionDeclaration` instance, while *body* is a - :class:`Block`. - """ - - self.ndim=ndim - self.upper = "queue_.submit([&](sycl::handler &h) {"+"\n h.parallel_for(range_, [=](sycl::nd_item<{}> item)".format(self.ndim) - self.body = body - self.lower=");\n }).wait();" + def __init__(self, body,ndim,handler,nd_item): + self.body =SYCLQueueSubmit(Block([SYCLparallel_for(body,ndim,handler,nd_item)]),handler) def generate(self): - yield self.upper yield from self.body.generate() - yield self.lower mapper_method = "map_function_body" + +class SYCLparallel_for(Generable): + def __init__(self, body,ndim,handler,nd_item): + self.handler=handler + self.args=["range_",FunctionBody(Lamda("=",[Value(f"sycl::nd_item<{ndim}>",nd_item)]),body)] + + def generate(self): + yield "{}.parallel_for({});".format(self.handler,", ".join(str(ad) for ad in self.args)) + + mapper_method = "map_SYCL_parallel_for" +class SYCLQueueSubmit(Generable): + def __init__(self, body,handler): + self.args=[FunctionBody(Lamda("&",[Value("sycl::handler &", handler)]),body)] + + def generate(self): + yield "queue_.submit({}).wait();".format(", ".join(str(ad) for ad in self.args)) + + mapper_method = "map_SYCL_queue_submit" \ No newline at end of file From f9e20f3d793027d37346f5536f50dbd4ab2cdaa7 Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Thu, 9 Mar 2023 00:51:44 +0530 Subject: [PATCH 4/8] Add SYCL local, global, constant --- cgen/sycl.py | 104 +++++++++++---------------------------------------- 1 file changed, 21 insertions(+), 83 deletions(-) diff --git a/cgen/sycl.py b/cgen/sycl.py index 7f66c0d..973d4bc 100644 --- a/cgen/sycl.py +++ b/cgen/sycl.py @@ -25,9 +25,6 @@ DeclSpecifier, \ NestedDeclarator, \ Value,\ - Pointer,\ - Extern,\ - NamespaceQualifier,\ Generable,\ FunctionBody,\ Lamda,\ @@ -71,43 +68,48 @@ def __init__(self, subdecl,ndim): subdecl.arg_decls.append(Value("sycl::queue","queue_")) subdecl.arg_decls.append(Value("sycl::nd_range<{}> ".format(self.ndim),"range_")) DeclSpecifier.__init__(self, subdecl,"") - - mapper_method = "map_cl_kernel" + mapper_method = "map_sycl_kernel" # }}} # {{{ kernel args -class CLConstant(DeclSpecifier): +class SYCLConstant(DeclSpecifier): def __init__(self, subdecl): - DeclSpecifier.__init__(self, subdecl, "__constant") + self.subdecl=subdecl - mapper_method = "map_cl_constant" + def get_decl_pair(self): + sub_tp,sub_decl=self.subdecl.get_decl_pair() + return [f"sycl::constant_ptr<{sub_tp[0]}>"],sub_decl + + mapper_method = "map_sycl_constant" class SYCLLocal(NestedDeclarator): def __init__(self, subdecl): - print(subdecl) - self.type=subdecl.get_type() self.subdecl=subdecl def get_decl_pair(self): sub_tp,sub_decl=self.subdecl.get_decl_pair() - return [f"sycl::local_accessor<{self.type}>"],sub_decl + return [f"sycl::local_ptr<{sub_tp[0]}>"],sub_decl mapper_method = "map_sycl_local" -class CLGlobal(DeclSpecifier): +class SYCLGlobal(DeclSpecifier): def __init__(self, subdecl): - DeclSpecifier.__init__(self, subdecl, "__global") - - mapper_method = "map_cl_global" + self.subdecl=subdecl + def get_decl_pair(self): + sub_tp,sub_decl=self.subdecl.get_decl_pair() + return [f"sycl::global_ptr<{sub_tp[0]}>"],sub_decl + + mapper_method = "map_sycl_global" -class CLImage(Value): +# TODO sycl Image +class SYCLImage(Value): def __init__(self, dims, mode, name): if mode == "r": spec = "__read_only" @@ -124,32 +126,6 @@ def __init__(self, dims, mode, name): # {{{ function attributes - -class CLVecTypeHint(NestedDeclarator): - def __init__(self, subdecl, dtype=None, count=None, type_str=None): - if (dtype is None) != (count is None): - raise ValueError("dtype and count must always be " - "specified together") - - if (dtype is None and type_str is None) or \ - (dtype is not None and type_str is not None): - raise ValueError("exactly one of dtype and type_str must be specified") - - if type_str is None: - self.type_str = dtype_to_cltype(dtype)+str(count) - else: - self.type_str = type_str - - super().__init__(subdecl) - - def get_decl_pair(self): - sub_tp, sub_decl = self.subdecl.get_decl_pair() - return sub_tp, ("__attribute__ ((vec_type_hint({}))) {}".format( - sub_decl, self.type_str)) - - mapper_method = "map_cl_vec_type_hint" - - class _SYCLWorkGroupSizeDeclarator(Generable): def __init__(self, dim): @@ -160,53 +136,15 @@ def __init__(self, dim): def generate(self): yield self.decl - -class CLWorkGroupSizeHint(_SYCLWorkGroupSizeDeclarator): - """ - See Sec 6.7.2 of OpenCL 2.0 spec, Version V2.2-11. - """ - - mapper_method = "map_cl_workgroup_size_hint" - - class SYCLRequiredWorkGroupSize(_SYCLWorkGroupSizeDeclarator): - """ - See Sec 6.7.2 of OpenCL 2.0 spec, Version V2.2-11. - """ - mapper_method = "map_cl_required_wokgroup_size" + mapper_method = "map_SYCL_required_wokgroup_size" # }}} - +#TODO SYCL vector # {{{ vector PODs -class CLVectorPOD(Declarator): - def __init__(self, dtype, count, name): - self.dtype = np.dtype(dtype) - self.count = count - self.name = name - - def get_decl_pair(self): - return [dtype_to_cltype(self.dtype)+str(self.count)], self.name - - def struct_maker_code(self, name): - return name - - def struct_format(self): - return str(self.count)+self.dtype.char - - def alignment_requirement(self): - from struct import calcsize - return calcsize(self.struct_format()) - - def default_value(self): - return [0]*self.count - - mapper_method = "map_cl_vector_pod" - # }}} - -# vim: fdm=marker class SYCLAccessor(NestedDeclarator): def __init__(self, subdecl,handler,count=None): NestedDeclarator.__init__(self, subdecl) @@ -234,7 +172,7 @@ def __init__(self, body,ndim,handler,nd_item): def generate(self): yield from self.body.generate() - mapper_method = "map_function_body" + mapper_method = "map_SYCL_function_body" class SYCLparallel_for(Generable): def __init__(self, body,ndim,handler,nd_item): From d175fba85a142a158ee0767b00337b252da88368 Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Thu, 9 Mar 2023 14:31:18 +0530 Subject: [PATCH 5/8] Change SYCLKernel, SYCLBody functions --- cgen/sycl.py | 122 ++++++++++++++++++++++++++++++++------------------- 1 file changed, 77 insertions(+), 45 deletions(-) diff --git a/cgen/sycl.py b/cgen/sycl.py index 973d4bc..1ffedc0 100644 --- a/cgen/sycl.py +++ b/cgen/sycl.py @@ -20,15 +20,15 @@ import numpy as np -from cgen import \ - Declarator, \ - DeclSpecifier, \ - NestedDeclarator, \ - Value,\ - Generable,\ - FunctionBody,\ - Lamda,\ - Block +from cgen import ( + Block, + DeclSpecifier, + FunctionBody, + Generable, + Lamda, + NestedDeclarator, + Value, +) def dtype_to_cltype(dtype): @@ -62,52 +62,58 @@ def dtype_to_cltype(dtype): # {{{ kernel + class SYCLKernel(DeclSpecifier): - def __init__(self, subdecl,ndim): - self.ndim=ndim - subdecl.arg_decls.append(Value("sycl::queue","queue_")) - subdecl.arg_decls.append(Value("sycl::nd_range<{}> ".format(self.ndim),"range_")) - DeclSpecifier.__init__(self, subdecl,"") + def __init__(self, subdecl, ndim, ndrange, queue): + self.ndim = ndim + subdecl.arg_decls.append(Value("sycl::queue", queue)) + subdecl.arg_decls.append( + Value("sycl::nd_range<{}> ".format(self.ndim), ndrange) + ) + DeclSpecifier.__init__(self, subdecl, "") mapper_method = "map_sycl_kernel" + # }}} # {{{ kernel args + class SYCLConstant(DeclSpecifier): def __init__(self, subdecl): - self.subdecl=subdecl + self.subdecl = subdecl def get_decl_pair(self): - sub_tp,sub_decl=self.subdecl.get_decl_pair() - return [f"sycl::constant_ptr<{sub_tp[0]}>"],sub_decl - + sub_tp, sub_decl = self.subdecl.get_decl_pair() + return [f"sycl::constant_ptr<{sub_tp[0]}>"], sub_decl + mapper_method = "map_sycl_constant" class SYCLLocal(NestedDeclarator): def __init__(self, subdecl): - self.subdecl=subdecl + self.subdecl = subdecl def get_decl_pair(self): - sub_tp,sub_decl=self.subdecl.get_decl_pair() - return [f"sycl::local_ptr<{sub_tp[0]}>"],sub_decl - + sub_tp, sub_decl = self.subdecl.get_decl_pair() + return [f"sycl::local_ptr<{sub_tp[0]}>"], sub_decl + mapper_method = "map_sycl_local" class SYCLGlobal(DeclSpecifier): def __init__(self, subdecl): - self.subdecl=subdecl + self.subdecl = subdecl def get_decl_pair(self): - sub_tp,sub_decl=self.subdecl.get_decl_pair() - return [f"sycl::global_ptr<{sub_tp[0]}>"],sub_decl - + sub_tp, sub_decl = self.subdecl.get_decl_pair() + return [f"sycl::global_ptr<{sub_tp[0]}>"], sub_decl + mapper_method = "map_sycl_global" + # TODO sycl Image class SYCLImage(Value): def __init__(self, dims, mode, name): @@ -122,6 +128,7 @@ def __init__(self, dims, mode, name): mapper_method = "map_cl_image" + # }}} @@ -132,62 +139,87 @@ def __init__(self, dim): while len(dim) < 3: dim = dim + (1,) self.dim = dim - self.decl='extern "C" [[sycl::reqd_work_group_size({})]]'.format(", ".join(str(d) for d in self.dim)) + self.decl = 'extern "C" [[sycl::reqd_work_group_size({})]]'.format( + ", ".join(str(d) for d in self.dim) + ) def generate(self): yield self.decl + + class SYCLRequiredWorkGroupSize(_SYCLWorkGroupSizeDeclarator): mapper_method = "map_SYCL_required_wokgroup_size" + # }}} -#TODO SYCL vector +# TODO SYCL vector # {{{ vector PODs # }}} class SYCLAccessor(NestedDeclarator): - def __init__(self, subdecl,handler,count=None): + def __init__(self, subdecl, handler, count=None): NestedDeclarator.__init__(self, subdecl) - self.handler=handler - self.count=count + self.handler = handler + self.count = count sub_tp, sub_decl = subdecl.get_decl_pair() - self.type=sub_tp[0] - self.sub_decl=sub_decl + self.type = sub_tp[0] + self.sub_decl = sub_decl def get_decl_pair(self): if self.count is None: return [f"sycl::accessor<{ self.type}>"], f"{self.sub_decl}({self.handler})" else: - return [f"sycl::accessor<{ self.type}>"], f"{self.sub_decl}(sycl::range<1>({self.count}),{self.handler})" + return [ + f"sycl::accessor<{ self.type}>" + ], f"{self.sub_decl}(sycl::range<1>({self.count}),{self.handler})" def get_type(self): return self.type + mapper_method = "map_sycl_accessor" - + class SYCLBody(Generable): - def __init__(self, body,ndim,handler,nd_item): - self.body =SYCLQueueSubmit(Block([SYCLparallel_for(body,ndim,handler,nd_item)]),handler) + def __init__(self, body, ndim, handler, nd_item, queue, ndrange): + self.body = SYCLQueueSubmit( + Block([SYCLparallel_for(body, ndim, handler, nd_item, ndrange)]), + handler, + queue, + ) def generate(self): yield from self.body.generate() mapper_method = "map_SYCL_function_body" + class SYCLparallel_for(Generable): - def __init__(self, body,ndim,handler,nd_item): - self.handler=handler - self.args=["range_",FunctionBody(Lamda("=",[Value(f"sycl::nd_item<{ndim}>",nd_item)]),body)] + def __init__(self, body, ndim, handler, nd_item, ndrange): + self.handler = handler + self.args = [ + ndrange, + FunctionBody(Lamda("=", [Value(f"sycl::nd_item<{ndim}>", nd_item)]), body), + ] def generate(self): - yield "{}.parallel_for({});".format(self.handler,", ".join(str(ad) for ad in self.args)) + yield "{}.parallel_for({});".format( + self.handler, ", ".join(str(ad) for ad in self.args) + ) mapper_method = "map_SYCL_parallel_for" + + class SYCLQueueSubmit(Generable): - def __init__(self, body,handler): - self.args=[FunctionBody(Lamda("&",[Value("sycl::handler &", handler)]),body)] + def __init__(self, body, handler, queue): + self.queue = queue + self.args = [ + FunctionBody(Lamda("&", [Value("sycl::handler &", handler)]), body) + ] def generate(self): - yield "queue_.submit({}).wait();".format(", ".join(str(ad) for ad in self.args)) + yield "{}.submit({}).wait();".format( + self.queue, ", ".join(str(ad) for ad in self.args) + ) - mapper_method = "map_SYCL_queue_submit" \ No newline at end of file + mapper_method = "map_SYCL_queue_submit" From b978ff898fee355784c9b7b76d57732aeefee603 Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Thu, 23 Mar 2023 12:20:15 +0530 Subject: [PATCH 6/8] Remove kernel sync --- cgen/sycl.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cgen/sycl.py b/cgen/sycl.py index 1ffedc0..e64a024 100644 --- a/cgen/sycl.py +++ b/cgen/sycl.py @@ -218,7 +218,7 @@ def __init__(self, body, handler, queue): ] def generate(self): - yield "{}.submit({}).wait();".format( + yield "{}.submit({});".format( self.queue, ", ".join(str(ad) for ad in self.args) ) From 04c4df84214ec3f8f852f4582c5f229077d14d8f Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Thu, 6 Apr 2023 12:51:27 +0530 Subject: [PATCH 7/8] Minor changes --- cgen/__init__.py | 8 ++++++-- cgen/sycl.py | 10 ++++++---- 2 files changed, 12 insertions(+), 6 deletions(-) diff --git a/cgen/__init__.py b/cgen/__init__.py index 8b89f90..728e6a6 100644 --- a/cgen/__init__.py +++ b/cgen/__init__.py @@ -413,18 +413,22 @@ def struct_format(self): mapper_method = "map_function_declaration" # }}} + + class Lamda(Declarator): - def __init__(self, capture_clause ,arg_decls): + def __init__(self, capture_clause, arg_decls): self.capture_clause = capture_clause self.arg_decls = arg_decls def get_decl_pair(self): - return [f"[{self.capture_clause}]"], f"({'{}'.format(', '.join(ad.inline() for ad in self.arg_decls))})" + arg_decls = ", ".join(ad.inline() for ad in self.arg_decls) + return [f"[{self.capture_clause}]"], f"({arg_decls})" mapper_method = "map_lamda" # {{{ struct-like + class Struct(Declarator): """A structure declarator.""" diff --git a/cgen/sycl.py b/cgen/sycl.py index e64a024..8db778d 100644 --- a/cgen/sycl.py +++ b/cgen/sycl.py @@ -168,7 +168,8 @@ def __init__(self, subdecl, handler, count=None): def get_decl_pair(self): if self.count is None: - return [f"sycl::accessor<{ self.type}>"], f"{self.sub_decl}({self.handler})" + return [f"sycl::accessor<{ self.type}>" + ], f"{self.sub_decl}({self.handler})" else: return [ f"sycl::accessor<{ self.type}>" @@ -183,7 +184,7 @@ def get_type(self): class SYCLBody(Generable): def __init__(self, body, ndim, handler, nd_item, queue, ndrange): self.body = SYCLQueueSubmit( - Block([SYCLparallel_for(body, ndim, handler, nd_item, ndrange)]), + Block([SYCLParallelFor(body, ndim, handler, nd_item, ndrange)]), handler, queue, ) @@ -194,12 +195,13 @@ def generate(self): mapper_method = "map_SYCL_function_body" -class SYCLparallel_for(Generable): +class SYCLParallelFor(Generable): def __init__(self, body, ndim, handler, nd_item, ndrange): self.handler = handler self.args = [ ndrange, - FunctionBody(Lamda("=", [Value(f"sycl::nd_item<{ndim}>", nd_item)]), body), + FunctionBody(Lamda("=", + [Value(f"sycl::nd_item<{ndim}>", nd_item)]), body), ] def generate(self): From 2c4e8812ae3ea84e05bd182f5b230ee00953b17f Mon Sep 17 00:00:00 2001 From: pubuduprabhathiya Date: Thu, 11 May 2023 14:58:35 +0530 Subject: [PATCH 8/8] Fixed SYCL local variable issue --- cgen/__init__.py | 7 +++++++ cgen/sycl.py | 17 +++++++++++++++-- 2 files changed, 22 insertions(+), 2 deletions(-) diff --git a/cgen/__init__.py b/cgen/__init__.py index 728e6a6..83e353a 100644 --- a/cgen/__init__.py +++ b/cgen/__init__.py @@ -389,6 +389,13 @@ def alignment_requirement(self): def default_value(self): return self.count*[self.subdecl.default_value()] + def get_size(self): + return self.count + + def get_array_name(self): + sub_tp, sub_decl = self.subdecl.get_decl_pair() + return str(sub_decl) + mapper_method = "map_array_of" diff --git a/cgen/sycl.py b/cgen/sycl.py index 8db778d..e439105 100644 --- a/cgen/sycl.py +++ b/cgen/sycl.py @@ -28,6 +28,8 @@ Lamda, NestedDeclarator, Value, + Assign, + ArrayOf ) @@ -93,12 +95,23 @@ def get_decl_pair(self): class SYCLLocal(NestedDeclarator): - def __init__(self, subdecl): + def __init__(self, subdecl, item): + self.item = item self.subdecl = subdecl def get_decl_pair(self): sub_tp, sub_decl = self.subdecl.get_decl_pair() - return [f"sycl::local_ptr<{sub_tp[0]}>"], sub_decl + if isinstance(self.subdecl, ArrayOf): + return ["auto&"], Assign( + self.subdecl.get_array_name(), + "*sycl::ext::oneapi::group_local_memory" + f"<{sub_tp[0]}[{self.subdecl.get_size()}]>" + f"({self.item}.get_group())") + else: + return ["auto&"], Assign( + sub_decl, + f"*sycl::ext::oneapi::group_local_memory<{sub_tp[0]}>" + f"({self.item}.get_group())") mapper_method = "map_sycl_local"