From 957b2b7e56d444f2179bcad6e39ca9f2c8e2bc83 Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Mon, 25 Nov 2024 14:47:43 +0100
Subject: [PATCH 01/13] First draft to add support for sycl accessors: Added a
function to add sycl accessors from a ps.field. Extended the parrallel_for to
make it possible add more statements before the actual kernel call
---
src/pystencilssfg/extensions/sycl.py | 8 ++-
src/pystencilssfg/lang/cpp/sycl_accessor.py | 77 +++++++++++++++++++++
2 files changed, 83 insertions(+), 2 deletions(-)
create mode 100644 src/pystencilssfg/lang/cpp/sycl_accessor.py
diff --git a/src/pystencilssfg/extensions/sycl.py b/src/pystencilssfg/extensions/sycl.py
index 3cb0c1c..bb99101 100644
--- a/src/pystencilssfg/extensions/sycl.py
+++ b/src/pystencilssfg/extensions/sycl.py
@@ -6,6 +6,8 @@ import re
from pystencils.types import PsType, PsCustomType
from pystencils.enums import Target
+from pystencilssfg.composer.basic_composer import SequencerArg
+
from ..exceptions import SfgException
from ..context import SfgContext
from ..composer import (
@@ -13,6 +15,7 @@ from ..composer import (
SfgClassComposer,
SfgComposer,
SfgComposerMixIn,
+ make_sequence,
)
from ..ir.source_components import SfgKernelHandle, SfgHeaderInclude
from ..ir import (
@@ -56,12 +59,13 @@ class SyclHandler(AugExpr):
self._ctx = ctx
- def parallel_for(self, range: SfgVar | Sequence[int], kernel: SfgKernelHandle):
+ def parallel_for(self, range: SfgVar | Sequence[int], kernel: SfgKernelHandle, *, extras: Sequence[SequencerArg]=[]):
"""Generate a ``parallel_for`` kernel invocation using this command group handler.
Args:
range: Object, or tuple of integers, indicating the kernel's iteration range
kernel: Handle to the pystencils-kernel to be executed
+ extras: Statements that should be in the parallel_for but before the kernel call
"""
self._ctx.add_include(SfgHeaderInclude("sycl/sycl.hpp", system_header=True))
@@ -81,7 +85,7 @@ class SyclHandler(AugExpr):
id_param = list(filter(filter_id, kernel.scalar_parameters))[0]
- tree = SfgKernelCallNode(kernel)
+ tree = make_sequence(*extras, SfgKernelCallNode(kernel))
kernel_lambda = SfgLambda(("=",), (id_param,), tree, None)
return SyclKernelInvoke(self, SyclInvokeType.ParallelFor, range, kernel_lambda)
diff --git a/src/pystencilssfg/lang/cpp/sycl_accessor.py b/src/pystencilssfg/lang/cpp/sycl_accessor.py
new file mode 100644
index 0000000..c6cd011
--- /dev/null
+++ b/src/pystencilssfg/lang/cpp/sycl_accessor.py
@@ -0,0 +1,77 @@
+from ...lang import SrcField, IFieldExtraction
+from ...ir.source_components import SfgHeaderInclude
+
+from pystencils import Field
+from pystencils.types import (
+ PsType,
+ PsCustomType,
+)
+
+from pystencilssfg.lang.expressions import AugExpr
+
+
+class SyclAccessor(SrcField):
+ def __init__(
+ self,
+ T: PsType,
+ dimensions: int,
+ reference: bool = False,
+ ):
+ cpp_typestr = T.c_string()
+ if dimensions not in [1, 2, 3]:
+ raise ValueError("sycl accessors can only have dims 1, 2 or 3")
+ typestring = (
+ f"sycl::accessor< {cpp_typestr}, {dimensions} > {'&' if reference else ''}"
+ )
+ super().__init__(PsCustomType(typestring))
+
+ self._dim = dimensions
+
+ @property
+ def required_includes(self) -> set[SfgHeaderInclude]:
+ return {SfgHeaderInclude("sycl/sycl.hpp", system_header=True)}
+
+ def get_extraction(self) -> IFieldExtraction:
+ accessor = self
+
+ class Extraction(IFieldExtraction):
+ def ptr(self) -> AugExpr:
+ return AugExpr.format(
+ "{}.get_multi_ptr<sycl::access::decorated::no>().get()",
+ accessor,
+ )
+
+ def size(self, coordinate: int) -> AugExpr | None:
+ if coordinate > accessor._dim:
+ return None
+ else:
+ return AugExpr.format(
+ "{}.get_range().get({})", accessor, coordinate
+ )
+
+ def stride(self, coordinate: int) -> AugExpr | None:
+ if coordinate > accessor._dim:
+ return None
+ else:
+ if coordinate == accessor._dim - 1:
+ return AugExpr.format("1")
+ else:
+ exprs = []
+ args = []
+ for d in range(coordinate + 1, accessor._dim):
+ args.extend([accessor, d])
+ exprs.append("{}.get_range().get({})")
+ expr = " * ".join(exprs)
+ return AugExpr.format(expr, *args)
+
+ return Extraction()
+
+
+def sycl_accessor_ref(field: Field):
+ """Creates a `sycl::accessor &` for a given pystencils field."""
+
+ return SyclAccessor(
+ field.dtype,
+ field.spatial_dimensions,
+ reference=True,
+ ).var(field.name)
--
GitLab
From 670d03c26967a5cbca6a6b700095a3206d7bd44d Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Mon, 25 Nov 2024 15:07:09 +0100
Subject: [PATCH 02/13] formatting
---
src/pystencilssfg/extensions/sycl.py | 8 +++++++-
1 file changed, 7 insertions(+), 1 deletion(-)
diff --git a/src/pystencilssfg/extensions/sycl.py b/src/pystencilssfg/extensions/sycl.py
index bb99101..2fdf0b1 100644
--- a/src/pystencilssfg/extensions/sycl.py
+++ b/src/pystencilssfg/extensions/sycl.py
@@ -59,7 +59,13 @@ class SyclHandler(AugExpr):
self._ctx = ctx
- def parallel_for(self, range: SfgVar | Sequence[int], kernel: SfgKernelHandle, *, extras: Sequence[SequencerArg]=[]):
+ def parallel_for(
+ self,
+ range: SfgVar | Sequence[int],
+ kernel: SfgKernelHandle,
+ *,
+ extras: Sequence[SequencerArg] = [],
+ ):
"""Generate a ``parallel_for`` kernel invocation using this command group handler.
Args:
--
GitLab
From 9672cf5758b29d8793e5bdf63ee021086735dca8 Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Tue, 26 Nov 2024 16:00:12 +0100
Subject: [PATCH 03/13] reworked the api for parallel_for to look like the
sfg.function api and extended the sycl_accessor mapping to handle index
dimensions
---
.gitignore | 3 +-
integration/test_sycl_buffer.py | 34 +++++++++++++++
src/pystencilssfg/extensions/sycl.py | 48 ++++++++++++++-------
src/pystencilssfg/lang/cpp/sycl_accessor.py | 41 ++++++++++++------
4 files changed, 97 insertions(+), 29 deletions(-)
create mode 100644 integration/test_sycl_buffer.py
diff --git a/.gitignore b/.gitignore
index 65c4064..399dabe 100644
--- a/.gitignore
+++ b/.gitignore
@@ -9,6 +9,7 @@
# dev environment
**/.venv
+**/venv
# build artifacts
dist
@@ -22,4 +23,4 @@ htmlcov
coverage.xml
# mkdocs
-site
\ No newline at end of file
+site
diff --git a/integration/test_sycl_buffer.py b/integration/test_sycl_buffer.py
new file mode 100644
index 0000000..7839b9f
--- /dev/null
+++ b/integration/test_sycl_buffer.py
@@ -0,0 +1,34 @@
+from pystencils import Target, CreateKernelConfig, no_jit
+from lbmpy import create_lb_update_rule, LBMOptimisation
+from pystencilssfg import SourceFileGenerator, SfgConfiguration, SfgOutputMode
+from pystencilssfg.lang.cpp.sycl_accessor import sycl_accessor_ref
+import pystencilssfg.extensions.sycl as sycl
+from itertools import chain
+
+sfg_config = SfgConfiguration(
+ output_directory="out/test_sycl_buffer",
+ outer_namespace="gen_code",
+ impl_extension="ipp",
+ output_mode=SfgOutputMode.INLINE,
+)
+
+with SourceFileGenerator(sfg_config) as sfg:
+ sfg = sycl.SyclComposer(sfg)
+ gen_config = CreateKernelConfig(target=Target.SYCL, jit=no_jit)
+ opt = LBMOptimisation(field_layout="fzyx")
+ update = create_lb_update_rule(lbm_optimisation=opt)
+ kernel = sfg.kernels.create(update, "lbm_update", gen_config)
+
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(update.method.dim, "range")
+ mappings = [
+ sfg.map_field(field, sycl_accessor_ref(field))
+ for field in chain(update.free_fields, update.bound_fields)
+ ]
+
+ sfg.function("lb_update")(
+ cgh.parallel_for(rang)(
+ *mappings,
+ sfg.call(kernel),
+ ),
+ )
diff --git a/src/pystencilssfg/extensions/sycl.py b/src/pystencilssfg/extensions/sycl.py
index 2fdf0b1..2f3591c 100644
--- a/src/pystencilssfg/extensions/sycl.py
+++ b/src/pystencilssfg/extensions/sycl.py
@@ -62,24 +62,29 @@ class SyclHandler(AugExpr):
def parallel_for(
self,
range: SfgVar | Sequence[int],
- kernel: SfgKernelHandle,
- *,
- extras: Sequence[SequencerArg] = [],
):
"""Generate a ``parallel_for`` kernel invocation using this command group handler.
+ The syntax of this uses a chain of two calls to mimic C++ syntax:
+
+ .. code-block:: Python
+
+ sfg.parallel_for(range)(
+ # Body
+ )
+
+ The body is constructed via sequencing (see `make_sequence`).
Args:
range: Object, or tuple of integers, indicating the kernel's iteration range
- kernel: Handle to the pystencils-kernel to be executed
- extras: Statements that should be in the parallel_for but before the kernel call
"""
self._ctx.add_include(SfgHeaderInclude("sycl/sycl.hpp", system_header=True))
- kfunc = kernel.get_kernel_function()
- if kfunc.target != Target.SYCL:
- raise SfgException(
- f"Kernel given to `parallel_for` is no SYCL kernel: {kernel.kernel_name}"
- )
+ def check_kernel(kernel: SfgKernelHandle):
+ kfunc = kernel.get_kernel_function()
+ if kfunc.target != Target.SYCL:
+ raise SfgException(
+ f"Kernel given to `parallel_for` is no SYCL kernel: {kernel.kernel_name}"
+ )
id_regex = re.compile(r"sycl::(id|item|nd_item)<\s*[0-9]\s*>")
@@ -89,12 +94,25 @@ class SyclHandler(AugExpr):
and id_regex.search(param.dtype.c_string()) is not None
)
- id_param = list(filter(filter_id, kernel.scalar_parameters))[0]
-
- tree = make_sequence(*extras, SfgKernelCallNode(kernel))
+ def sequencer(*args: SequencerArg):
+ id_param = []
+ for arg in args:
+ if isinstance(arg, SfgKernelCallNode):
+ check_kernel(arg._kernel_handle)
+ id_param.append(list(filter(filter_id, arg._kernel_handle.scalar_parameters))[0])
+
+ if not all(item == id_param[0] for item in id_param):
+ raise ValueError(
+ "id_param should be the same for all kernels in parallel_for"
+ )
+ tree = make_sequence(*args)
+
+ kernel_lambda = SfgLambda(("=",), (id_param[0],), tree, None)
+ return SyclKernelInvoke(
+ self, SyclInvokeType.ParallelFor, range, kernel_lambda
+ )
- kernel_lambda = SfgLambda(("=",), (id_param,), tree, None)
- return SyclKernelInvoke(self, SyclInvokeType.ParallelFor, range, kernel_lambda)
+ return sequencer
class SyclGroup(AugExpr):
diff --git a/src/pystencilssfg/lang/cpp/sycl_accessor.py b/src/pystencilssfg/lang/cpp/sycl_accessor.py
index c6cd011..e8618a5 100644
--- a/src/pystencilssfg/lang/cpp/sycl_accessor.py
+++ b/src/pystencilssfg/lang/cpp/sycl_accessor.py
@@ -1,5 +1,7 @@
+import math
from ...lang import SrcField, IFieldExtraction
from ...ir.source_components import SfgHeaderInclude
+from typing import Sequence
from pystencils import Field
from pystencils.types import (
@@ -15,6 +17,7 @@ class SyclAccessor(SrcField):
self,
T: PsType,
dimensions: int,
+ index_shape: Sequence[int],
reference: bool = False,
):
cpp_typestr = T.c_string()
@@ -25,7 +28,11 @@ class SyclAccessor(SrcField):
)
super().__init__(PsCustomType(typestring))
- self._dim = dimensions
+ self._spatial_dimensions = dimensions
+ self._index_dimensions = len(index_shape)
+ self._index_shape = index_shape
+ self._index_size = math.prod(index_shape)
+ self._total_dimensions_ = self._spatial_dimensions + self._index_dimensions
@property
def required_includes(self) -> set[SfgHeaderInclude]:
@@ -42,7 +49,7 @@ class SyclAccessor(SrcField):
)
def size(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._dim:
+ if coordinate > accessor._spatial_dimensions:
return None
else:
return AugExpr.format(
@@ -50,28 +57,36 @@ class SyclAccessor(SrcField):
)
def stride(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._dim:
+ if coordinate > accessor._total_dimensions_:
return None
+ elif coordinate >= accessor._spatial_dimensions - 1:
+ start = (coordinate - accessor._spatial_dimensions) + 1
+ return AugExpr.format(
+ "{}", math.prod(accessor._index_shape[start:])
+ )
else:
- if coordinate == accessor._dim - 1:
- return AugExpr.format("1")
- else:
- exprs = []
- args = []
- for d in range(coordinate + 1, accessor._dim):
- args.extend([accessor, d])
- exprs.append("{}.get_range().get({})")
- expr = " * ".join(exprs)
- return AugExpr.format(expr, *args)
+ exprs = []
+ args = []
+ for d in range(coordinate + 1, accessor._spatial_dimensions):
+ args.extend([accessor, d])
+ exprs.append("{}.get_range().get({})")
+ expr = " * ".join(exprs)
+ expr += " * {}"
+ return AugExpr.format(expr, *args, accessor._index_size)
return Extraction()
def sycl_accessor_ref(field: Field):
"""Creates a `sycl::accessor &` for a given pystencils field."""
+ # Sycl accesors allow only at max 3 dimensions:
+ # So also mapping the index dimens to the sycl accesor we only can have 2D LBM stuff
+ # In principle it would be possible to map it to something like sycl::buffer<std::array<double, 19>, 3>
+ # but then would need to generate kernels that have sycl accessors as arguments
return SyclAccessor(
field.dtype,
field.spatial_dimensions,
+ field.index_shape,
reference=True,
).var(field.name)
--
GitLab
From ef144a22de027eff64dc25e7180656c17dec2053 Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Thu, 28 Nov 2024 13:19:15 +0100
Subject: [PATCH 04/13] All dimensions of a field are now mapped to the
accessor dimension, this allows only only 3 dims at maximum (spatial and
index dim ) add some tests
---
src/pystencilssfg/lang/cpp/sycl_accessor.py | 37 ++++++-------------
tests/lang/test_sycl.py | 41 +++++++++++++++++++++
2 files changed, 53 insertions(+), 25 deletions(-)
create mode 100644 tests/lang/test_sycl.py
diff --git a/src/pystencilssfg/lang/cpp/sycl_accessor.py b/src/pystencilssfg/lang/cpp/sycl_accessor.py
index e8618a5..f704477 100644
--- a/src/pystencilssfg/lang/cpp/sycl_accessor.py
+++ b/src/pystencilssfg/lang/cpp/sycl_accessor.py
@@ -1,7 +1,5 @@
-import math
from ...lang import SrcField, IFieldExtraction
from ...ir.source_components import SfgHeaderInclude
-from typing import Sequence
from pystencils import Field
from pystencils.types import (
@@ -17,22 +15,17 @@ class SyclAccessor(SrcField):
self,
T: PsType,
dimensions: int,
- index_shape: Sequence[int],
reference: bool = False,
):
cpp_typestr = T.c_string()
- if dimensions not in [1, 2, 3]:
+ if 3 < dimensions:
raise ValueError("sycl accessors can only have dims 1, 2 or 3")
typestring = (
f"sycl::accessor< {cpp_typestr}, {dimensions} > {'&' if reference else ''}"
)
super().__init__(PsCustomType(typestring))
-
- self._spatial_dimensions = dimensions
- self._index_dimensions = len(index_shape)
- self._index_shape = index_shape
- self._index_size = math.prod(index_shape)
- self._total_dimensions_ = self._spatial_dimensions + self._index_dimensions
+ self._dim = dimensions
+ self._inner_stride = 1
@property
def required_includes(self) -> set[SfgHeaderInclude]:
@@ -49,7 +42,7 @@ class SyclAccessor(SrcField):
)
def size(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._spatial_dimensions:
+ if coordinate > accessor._dim:
return None
else:
return AugExpr.format(
@@ -57,36 +50,30 @@ class SyclAccessor(SrcField):
)
def stride(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._total_dimensions_:
+ if coordinate > accessor._dim:
return None
- elif coordinate >= accessor._spatial_dimensions - 1:
- start = (coordinate - accessor._spatial_dimensions) + 1
- return AugExpr.format(
- "{}", math.prod(accessor._index_shape[start:])
- )
+ elif coordinate == accessor._dim - 1:
+ return AugExpr.format("{}", accessor._inner_stride)
else:
exprs = []
args = []
- for d in range(coordinate + 1, accessor._spatial_dimensions):
+ for d in range(coordinate + 1, accessor._dim):
args.extend([accessor, d])
exprs.append("{}.get_range().get({})")
expr = " * ".join(exprs)
expr += " * {}"
- return AugExpr.format(expr, *args, accessor._index_size)
+ return AugExpr.format(expr, *args, accessor._inner_stride)
return Extraction()
def sycl_accessor_ref(field: Field):
"""Creates a `sycl::accessor &` for a given pystencils field."""
- # Sycl accesors allow only at max 3 dimensions:
- # So also mapping the index dimens to the sycl accesor we only can have 2D LBM stuff
- # In principle it would be possible to map it to something like sycl::buffer<std::array<double, 19>, 3>
- # but then would need to generate kernels that have sycl accessors as arguments
+ # Sycl Accessor do not expose information about strides, so the linearization is like here
+ # https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_multi_dimensional_objects_and_linearization
return SyclAccessor(
field.dtype,
- field.spatial_dimensions,
- field.index_shape,
+ field.spatial_dimensions + field.index_dimensions,
reference=True,
).var(field.name)
diff --git a/tests/lang/test_sycl.py b/tests/lang/test_sycl.py
new file mode 100644
index 0000000..88f958f
--- /dev/null
+++ b/tests/lang/test_sycl.py
@@ -0,0 +1,41 @@
+import pytest
+
+
+from pystencilssfg.lang.cpp.sycl_accessor import sycl_accessor_ref
+import pystencils as ps
+
+
+@pytest.mark.parametrize("data_type", ["double", "float"])
+@pytest.mark.parametrize("dim", [1, 2, 3])
+def test_spatial_field(data_type, dim):
+ f = ps.fields(f"f:{data_type}[{dim}D]")
+ ref = sycl_accessor_ref(f)
+ assert (
+ f"sycl::accessor< {data_type}, {dim} > &" in ref.as_variable().name_and_type()
+ )
+
+
+@pytest.mark.parametrize("data_type", ["double", "float"])
+def test_to_large_dim(data_type):
+ dim = 4
+ f = ps.fields(f"f:{data_type}[{dim}D]")
+ with pytest.raises(ValueError):
+ sycl_accessor_ref(f)
+
+
+@pytest.mark.parametrize("data_type", ["double", "float"])
+@pytest.mark.parametrize("spatial_dim", [1, 2, 3])
+@pytest.mark.parametrize("index_dims", [1, 2, 3])
+def test_index_field(data_type, spatial_dim, index_dims):
+ index_shape = ("19",) * index_dims
+ total_dims = spatial_dim + index_dims
+ f = ps.fields(f"f({', '.join(index_shape)}):{data_type}[{spatial_dim}D]")
+ if total_dims <= 3:
+ ref = sycl_accessor_ref(f)
+ assert (
+ f"sycl::accessor< {data_type}, {total_dims} > &"
+ in ref.as_variable().name_and_type()
+ )
+ else:
+ with pytest.raises(ValueError):
+ sycl_accessor_ref(f)
--
GitLab
From d862626f27396de12643ddfdbbd3316217b98845 Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Thu, 28 Nov 2024 13:54:14 +0100
Subject: [PATCH 05/13] added a test for the parallel_for
---
tests/extensions/test_sycl.py | 64 +++++++++++++++++++
.../{test_sycl.py => test_sycl_accessor.py} | 0
2 files changed, 64 insertions(+)
create mode 100644 tests/extensions/test_sycl.py
rename tests/lang/{test_sycl.py => test_sycl_accessor.py} (100%)
diff --git a/tests/extensions/test_sycl.py b/tests/extensions/test_sycl.py
new file mode 100644
index 0000000..db99278
--- /dev/null
+++ b/tests/extensions/test_sycl.py
@@ -0,0 +1,64 @@
+import pytest
+from pystencilssfg import SourceFileGenerator
+import pystencilssfg.extensions.sycl as sycl
+import pystencils as ps
+from pystencilssfg import SfgContext
+
+
+def test_parallel_for_1_kernels():
+ sfg = sycl.SyclComposer(SfgContext())
+ data_type = "double"
+ dim = 2
+ f, g, h, i = ps.fields(f"f,g,h,i:{data_type}[{dim}D]")
+ assignement_1 = ps.Assignment(f.center(), g.center())
+ assignement_2 = ps.Assignment(h.center(), i.center())
+
+ config = ps.CreateKernelConfig(target=ps.Target.SYCL)
+ kernel_1 = sfg.kernels.create(assignement_1, "kernel_1", config)
+ kernel_2 = sfg.kernels.create(assignement_2, "kernel_2", config)
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(dim, "range")
+ cgh.parallel_for(rang)(
+ sfg.call(kernel_1),
+ sfg.call(kernel_2),
+ )
+
+
+def test_parallel_for_2_kernels():
+ sfg = sycl.SyclComposer(SfgContext())
+ data_type = "double"
+ dim = 2
+ f, g, h, i = ps.fields(f"f,g,h,i:{data_type}[{dim}D]")
+ assignement_1 = ps.Assignment(f.center(), g.center())
+ assignement_2 = ps.Assignment(h.center(), i.center())
+
+ config = ps.CreateKernelConfig(target=ps.Target.SYCL)
+ kernel_1 = sfg.kernels.create(assignement_1, "kernel_1", config)
+ kernel_2 = sfg.kernels.create(assignement_2, "kernel_2", config)
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(dim, "range")
+ cgh.parallel_for(rang)(
+ sfg.call(kernel_1),
+ sfg.call(kernel_2),
+ )
+
+
+def test_parallel_for_2_kernels_fail():
+ sfg = sycl.SyclComposer(SfgContext())
+ data_type = "double"
+ dim = 2
+ f, g = ps.fields(f"f,g:{data_type}[{dim}D]")
+ h, i = ps.fields(f"h,i:{data_type}[{dim-1}D]")
+ assignement_1 = ps.Assignment(f.center(), g.center())
+ assignement_2 = ps.Assignment(h.center(), i.center())
+
+ config = ps.CreateKernelConfig(target=ps.Target.SYCL)
+ kernel_1 = sfg.kernels.create(assignement_1, "kernel_1", config)
+ kernel_2 = sfg.kernels.create(assignement_2, "kernel_2", config)
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(dim, "range")
+ with pytest.raises(ValueError):
+ cgh.parallel_for(rang)(
+ sfg.call(kernel_1),
+ sfg.call(kernel_2),
+ )
diff --git a/tests/lang/test_sycl.py b/tests/lang/test_sycl_accessor.py
similarity index 100%
rename from tests/lang/test_sycl.py
rename to tests/lang/test_sycl_accessor.py
--
GitLab
From 218c3334d922785836e40cd033172ec0e55a1891 Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Mon, 25 Nov 2024 14:47:43 +0100
Subject: [PATCH 06/13] First draft to add support for sycl accessors: Added a
function to add sycl accessors from a ps.field. Extended the parrallel_for to
make it possible add more statements before the actual kernel call
---
src/pystencilssfg/extensions/sycl.py | 8 ++-
src/pystencilssfg/lang/cpp/sycl_accessor.py | 77 +++++++++++++++++++++
2 files changed, 83 insertions(+), 2 deletions(-)
create mode 100644 src/pystencilssfg/lang/cpp/sycl_accessor.py
diff --git a/src/pystencilssfg/extensions/sycl.py b/src/pystencilssfg/extensions/sycl.py
index 3cb0c1c..bb99101 100644
--- a/src/pystencilssfg/extensions/sycl.py
+++ b/src/pystencilssfg/extensions/sycl.py
@@ -6,6 +6,8 @@ import re
from pystencils.types import PsType, PsCustomType
from pystencils.enums import Target
+from pystencilssfg.composer.basic_composer import SequencerArg
+
from ..exceptions import SfgException
from ..context import SfgContext
from ..composer import (
@@ -13,6 +15,7 @@ from ..composer import (
SfgClassComposer,
SfgComposer,
SfgComposerMixIn,
+ make_sequence,
)
from ..ir.source_components import SfgKernelHandle, SfgHeaderInclude
from ..ir import (
@@ -56,12 +59,13 @@ class SyclHandler(AugExpr):
self._ctx = ctx
- def parallel_for(self, range: SfgVar | Sequence[int], kernel: SfgKernelHandle):
+ def parallel_for(self, range: SfgVar | Sequence[int], kernel: SfgKernelHandle, *, extras: Sequence[SequencerArg]=[]):
"""Generate a ``parallel_for`` kernel invocation using this command group handler.
Args:
range: Object, or tuple of integers, indicating the kernel's iteration range
kernel: Handle to the pystencils-kernel to be executed
+ extras: Statements that should be in the parallel_for but before the kernel call
"""
self._ctx.add_include(SfgHeaderInclude("sycl/sycl.hpp", system_header=True))
@@ -81,7 +85,7 @@ class SyclHandler(AugExpr):
id_param = list(filter(filter_id, kernel.scalar_parameters))[0]
- tree = SfgKernelCallNode(kernel)
+ tree = make_sequence(*extras, SfgKernelCallNode(kernel))
kernel_lambda = SfgLambda(("=",), (id_param,), tree, None)
return SyclKernelInvoke(self, SyclInvokeType.ParallelFor, range, kernel_lambda)
diff --git a/src/pystencilssfg/lang/cpp/sycl_accessor.py b/src/pystencilssfg/lang/cpp/sycl_accessor.py
new file mode 100644
index 0000000..c6cd011
--- /dev/null
+++ b/src/pystencilssfg/lang/cpp/sycl_accessor.py
@@ -0,0 +1,77 @@
+from ...lang import SrcField, IFieldExtraction
+from ...ir.source_components import SfgHeaderInclude
+
+from pystencils import Field
+from pystencils.types import (
+ PsType,
+ PsCustomType,
+)
+
+from pystencilssfg.lang.expressions import AugExpr
+
+
+class SyclAccessor(SrcField):
+ def __init__(
+ self,
+ T: PsType,
+ dimensions: int,
+ reference: bool = False,
+ ):
+ cpp_typestr = T.c_string()
+ if dimensions not in [1, 2, 3]:
+ raise ValueError("sycl accessors can only have dims 1, 2 or 3")
+ typestring = (
+ f"sycl::accessor< {cpp_typestr}, {dimensions} > {'&' if reference else ''}"
+ )
+ super().__init__(PsCustomType(typestring))
+
+ self._dim = dimensions
+
+ @property
+ def required_includes(self) -> set[SfgHeaderInclude]:
+ return {SfgHeaderInclude("sycl/sycl.hpp", system_header=True)}
+
+ def get_extraction(self) -> IFieldExtraction:
+ accessor = self
+
+ class Extraction(IFieldExtraction):
+ def ptr(self) -> AugExpr:
+ return AugExpr.format(
+ "{}.get_multi_ptr<sycl::access::decorated::no>().get()",
+ accessor,
+ )
+
+ def size(self, coordinate: int) -> AugExpr | None:
+ if coordinate > accessor._dim:
+ return None
+ else:
+ return AugExpr.format(
+ "{}.get_range().get({})", accessor, coordinate
+ )
+
+ def stride(self, coordinate: int) -> AugExpr | None:
+ if coordinate > accessor._dim:
+ return None
+ else:
+ if coordinate == accessor._dim - 1:
+ return AugExpr.format("1")
+ else:
+ exprs = []
+ args = []
+ for d in range(coordinate + 1, accessor._dim):
+ args.extend([accessor, d])
+ exprs.append("{}.get_range().get({})")
+ expr = " * ".join(exprs)
+ return AugExpr.format(expr, *args)
+
+ return Extraction()
+
+
+def sycl_accessor_ref(field: Field):
+ """Creates a `sycl::accessor &` for a given pystencils field."""
+
+ return SyclAccessor(
+ field.dtype,
+ field.spatial_dimensions,
+ reference=True,
+ ).var(field.name)
--
GitLab
From 188308e1ea63bb84a1f094723a907dc7300fb941 Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Mon, 25 Nov 2024 15:07:09 +0100
Subject: [PATCH 07/13] formatting
---
src/pystencilssfg/extensions/sycl.py | 8 +++++++-
1 file changed, 7 insertions(+), 1 deletion(-)
diff --git a/src/pystencilssfg/extensions/sycl.py b/src/pystencilssfg/extensions/sycl.py
index bb99101..2fdf0b1 100644
--- a/src/pystencilssfg/extensions/sycl.py
+++ b/src/pystencilssfg/extensions/sycl.py
@@ -59,7 +59,13 @@ class SyclHandler(AugExpr):
self._ctx = ctx
- def parallel_for(self, range: SfgVar | Sequence[int], kernel: SfgKernelHandle, *, extras: Sequence[SequencerArg]=[]):
+ def parallel_for(
+ self,
+ range: SfgVar | Sequence[int],
+ kernel: SfgKernelHandle,
+ *,
+ extras: Sequence[SequencerArg] = [],
+ ):
"""Generate a ``parallel_for`` kernel invocation using this command group handler.
Args:
--
GitLab
From 26dc9e70e2dba9290f6916da96369215d7c84b5c Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Tue, 26 Nov 2024 16:00:12 +0100
Subject: [PATCH 08/13] reworked the api for parallel_for to look like the
sfg.function api and extended the sycl_accessor mapping to handle index
dimensions
---
.gitignore | 3 +-
integration/test_sycl_buffer.py | 34 +++++++++++++++
src/pystencilssfg/extensions/sycl.py | 48 ++++++++++++++-------
src/pystencilssfg/lang/cpp/sycl_accessor.py | 41 ++++++++++++------
4 files changed, 97 insertions(+), 29 deletions(-)
create mode 100644 integration/test_sycl_buffer.py
diff --git a/.gitignore b/.gitignore
index 65c4064..399dabe 100644
--- a/.gitignore
+++ b/.gitignore
@@ -9,6 +9,7 @@
# dev environment
**/.venv
+**/venv
# build artifacts
dist
@@ -22,4 +23,4 @@ htmlcov
coverage.xml
# mkdocs
-site
\ No newline at end of file
+site
diff --git a/integration/test_sycl_buffer.py b/integration/test_sycl_buffer.py
new file mode 100644
index 0000000..7839b9f
--- /dev/null
+++ b/integration/test_sycl_buffer.py
@@ -0,0 +1,34 @@
+from pystencils import Target, CreateKernelConfig, no_jit
+from lbmpy import create_lb_update_rule, LBMOptimisation
+from pystencilssfg import SourceFileGenerator, SfgConfiguration, SfgOutputMode
+from pystencilssfg.lang.cpp.sycl_accessor import sycl_accessor_ref
+import pystencilssfg.extensions.sycl as sycl
+from itertools import chain
+
+sfg_config = SfgConfiguration(
+ output_directory="out/test_sycl_buffer",
+ outer_namespace="gen_code",
+ impl_extension="ipp",
+ output_mode=SfgOutputMode.INLINE,
+)
+
+with SourceFileGenerator(sfg_config) as sfg:
+ sfg = sycl.SyclComposer(sfg)
+ gen_config = CreateKernelConfig(target=Target.SYCL, jit=no_jit)
+ opt = LBMOptimisation(field_layout="fzyx")
+ update = create_lb_update_rule(lbm_optimisation=opt)
+ kernel = sfg.kernels.create(update, "lbm_update", gen_config)
+
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(update.method.dim, "range")
+ mappings = [
+ sfg.map_field(field, sycl_accessor_ref(field))
+ for field in chain(update.free_fields, update.bound_fields)
+ ]
+
+ sfg.function("lb_update")(
+ cgh.parallel_for(rang)(
+ *mappings,
+ sfg.call(kernel),
+ ),
+ )
diff --git a/src/pystencilssfg/extensions/sycl.py b/src/pystencilssfg/extensions/sycl.py
index 2fdf0b1..2f3591c 100644
--- a/src/pystencilssfg/extensions/sycl.py
+++ b/src/pystencilssfg/extensions/sycl.py
@@ -62,24 +62,29 @@ class SyclHandler(AugExpr):
def parallel_for(
self,
range: SfgVar | Sequence[int],
- kernel: SfgKernelHandle,
- *,
- extras: Sequence[SequencerArg] = [],
):
"""Generate a ``parallel_for`` kernel invocation using this command group handler.
+ The syntax of this uses a chain of two calls to mimic C++ syntax:
+
+ .. code-block:: Python
+
+ sfg.parallel_for(range)(
+ # Body
+ )
+
+ The body is constructed via sequencing (see `make_sequence`).
Args:
range: Object, or tuple of integers, indicating the kernel's iteration range
- kernel: Handle to the pystencils-kernel to be executed
- extras: Statements that should be in the parallel_for but before the kernel call
"""
self._ctx.add_include(SfgHeaderInclude("sycl/sycl.hpp", system_header=True))
- kfunc = kernel.get_kernel_function()
- if kfunc.target != Target.SYCL:
- raise SfgException(
- f"Kernel given to `parallel_for` is no SYCL kernel: {kernel.kernel_name}"
- )
+ def check_kernel(kernel: SfgKernelHandle):
+ kfunc = kernel.get_kernel_function()
+ if kfunc.target != Target.SYCL:
+ raise SfgException(
+ f"Kernel given to `parallel_for` is no SYCL kernel: {kernel.kernel_name}"
+ )
id_regex = re.compile(r"sycl::(id|item|nd_item)<\s*[0-9]\s*>")
@@ -89,12 +94,25 @@ class SyclHandler(AugExpr):
and id_regex.search(param.dtype.c_string()) is not None
)
- id_param = list(filter(filter_id, kernel.scalar_parameters))[0]
-
- tree = make_sequence(*extras, SfgKernelCallNode(kernel))
+ def sequencer(*args: SequencerArg):
+ id_param = []
+ for arg in args:
+ if isinstance(arg, SfgKernelCallNode):
+ check_kernel(arg._kernel_handle)
+ id_param.append(list(filter(filter_id, arg._kernel_handle.scalar_parameters))[0])
+
+ if not all(item == id_param[0] for item in id_param):
+ raise ValueError(
+ "id_param should be the same for all kernels in parallel_for"
+ )
+ tree = make_sequence(*args)
+
+ kernel_lambda = SfgLambda(("=",), (id_param[0],), tree, None)
+ return SyclKernelInvoke(
+ self, SyclInvokeType.ParallelFor, range, kernel_lambda
+ )
- kernel_lambda = SfgLambda(("=",), (id_param,), tree, None)
- return SyclKernelInvoke(self, SyclInvokeType.ParallelFor, range, kernel_lambda)
+ return sequencer
class SyclGroup(AugExpr):
diff --git a/src/pystencilssfg/lang/cpp/sycl_accessor.py b/src/pystencilssfg/lang/cpp/sycl_accessor.py
index c6cd011..e8618a5 100644
--- a/src/pystencilssfg/lang/cpp/sycl_accessor.py
+++ b/src/pystencilssfg/lang/cpp/sycl_accessor.py
@@ -1,5 +1,7 @@
+import math
from ...lang import SrcField, IFieldExtraction
from ...ir.source_components import SfgHeaderInclude
+from typing import Sequence
from pystencils import Field
from pystencils.types import (
@@ -15,6 +17,7 @@ class SyclAccessor(SrcField):
self,
T: PsType,
dimensions: int,
+ index_shape: Sequence[int],
reference: bool = False,
):
cpp_typestr = T.c_string()
@@ -25,7 +28,11 @@ class SyclAccessor(SrcField):
)
super().__init__(PsCustomType(typestring))
- self._dim = dimensions
+ self._spatial_dimensions = dimensions
+ self._index_dimensions = len(index_shape)
+ self._index_shape = index_shape
+ self._index_size = math.prod(index_shape)
+ self._total_dimensions_ = self._spatial_dimensions + self._index_dimensions
@property
def required_includes(self) -> set[SfgHeaderInclude]:
@@ -42,7 +49,7 @@ class SyclAccessor(SrcField):
)
def size(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._dim:
+ if coordinate > accessor._spatial_dimensions:
return None
else:
return AugExpr.format(
@@ -50,28 +57,36 @@ class SyclAccessor(SrcField):
)
def stride(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._dim:
+ if coordinate > accessor._total_dimensions_:
return None
+ elif coordinate >= accessor._spatial_dimensions - 1:
+ start = (coordinate - accessor._spatial_dimensions) + 1
+ return AugExpr.format(
+ "{}", math.prod(accessor._index_shape[start:])
+ )
else:
- if coordinate == accessor._dim - 1:
- return AugExpr.format("1")
- else:
- exprs = []
- args = []
- for d in range(coordinate + 1, accessor._dim):
- args.extend([accessor, d])
- exprs.append("{}.get_range().get({})")
- expr = " * ".join(exprs)
- return AugExpr.format(expr, *args)
+ exprs = []
+ args = []
+ for d in range(coordinate + 1, accessor._spatial_dimensions):
+ args.extend([accessor, d])
+ exprs.append("{}.get_range().get({})")
+ expr = " * ".join(exprs)
+ expr += " * {}"
+ return AugExpr.format(expr, *args, accessor._index_size)
return Extraction()
def sycl_accessor_ref(field: Field):
"""Creates a `sycl::accessor &` for a given pystencils field."""
+ # Sycl accesors allow only at max 3 dimensions:
+ # So also mapping the index dimens to the sycl accesor we only can have 2D LBM stuff
+ # In principle it would be possible to map it to something like sycl::buffer<std::array<double, 19>, 3>
+ # but then would need to generate kernels that have sycl accessors as arguments
return SyclAccessor(
field.dtype,
field.spatial_dimensions,
+ field.index_shape,
reference=True,
).var(field.name)
--
GitLab
From 764f56e4147c28cf4f7a23a5dd29eddc854150e2 Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Thu, 28 Nov 2024 13:19:15 +0100
Subject: [PATCH 09/13] All dimensions of a field are now mapped to the
accessor dimension, this allows only only 3 dims at maximum (spatial and
index dim ) add some tests
---
src/pystencilssfg/lang/cpp/sycl_accessor.py | 37 ++++++-------------
tests/lang/test_sycl.py | 41 +++++++++++++++++++++
2 files changed, 53 insertions(+), 25 deletions(-)
create mode 100644 tests/lang/test_sycl.py
diff --git a/src/pystencilssfg/lang/cpp/sycl_accessor.py b/src/pystencilssfg/lang/cpp/sycl_accessor.py
index e8618a5..f704477 100644
--- a/src/pystencilssfg/lang/cpp/sycl_accessor.py
+++ b/src/pystencilssfg/lang/cpp/sycl_accessor.py
@@ -1,7 +1,5 @@
-import math
from ...lang import SrcField, IFieldExtraction
from ...ir.source_components import SfgHeaderInclude
-from typing import Sequence
from pystencils import Field
from pystencils.types import (
@@ -17,22 +15,17 @@ class SyclAccessor(SrcField):
self,
T: PsType,
dimensions: int,
- index_shape: Sequence[int],
reference: bool = False,
):
cpp_typestr = T.c_string()
- if dimensions not in [1, 2, 3]:
+ if 3 < dimensions:
raise ValueError("sycl accessors can only have dims 1, 2 or 3")
typestring = (
f"sycl::accessor< {cpp_typestr}, {dimensions} > {'&' if reference else ''}"
)
super().__init__(PsCustomType(typestring))
-
- self._spatial_dimensions = dimensions
- self._index_dimensions = len(index_shape)
- self._index_shape = index_shape
- self._index_size = math.prod(index_shape)
- self._total_dimensions_ = self._spatial_dimensions + self._index_dimensions
+ self._dim = dimensions
+ self._inner_stride = 1
@property
def required_includes(self) -> set[SfgHeaderInclude]:
@@ -49,7 +42,7 @@ class SyclAccessor(SrcField):
)
def size(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._spatial_dimensions:
+ if coordinate > accessor._dim:
return None
else:
return AugExpr.format(
@@ -57,36 +50,30 @@ class SyclAccessor(SrcField):
)
def stride(self, coordinate: int) -> AugExpr | None:
- if coordinate > accessor._total_dimensions_:
+ if coordinate > accessor._dim:
return None
- elif coordinate >= accessor._spatial_dimensions - 1:
- start = (coordinate - accessor._spatial_dimensions) + 1
- return AugExpr.format(
- "{}", math.prod(accessor._index_shape[start:])
- )
+ elif coordinate == accessor._dim - 1:
+ return AugExpr.format("{}", accessor._inner_stride)
else:
exprs = []
args = []
- for d in range(coordinate + 1, accessor._spatial_dimensions):
+ for d in range(coordinate + 1, accessor._dim):
args.extend([accessor, d])
exprs.append("{}.get_range().get({})")
expr = " * ".join(exprs)
expr += " * {}"
- return AugExpr.format(expr, *args, accessor._index_size)
+ return AugExpr.format(expr, *args, accessor._inner_stride)
return Extraction()
def sycl_accessor_ref(field: Field):
"""Creates a `sycl::accessor &` for a given pystencils field."""
- # Sycl accesors allow only at max 3 dimensions:
- # So also mapping the index dimens to the sycl accesor we only can have 2D LBM stuff
- # In principle it would be possible to map it to something like sycl::buffer<std::array<double, 19>, 3>
- # but then would need to generate kernels that have sycl accessors as arguments
+ # Sycl Accessor do not expose information about strides, so the linearization is like here
+ # https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_multi_dimensional_objects_and_linearization
return SyclAccessor(
field.dtype,
- field.spatial_dimensions,
- field.index_shape,
+ field.spatial_dimensions + field.index_dimensions,
reference=True,
).var(field.name)
diff --git a/tests/lang/test_sycl.py b/tests/lang/test_sycl.py
new file mode 100644
index 0000000..88f958f
--- /dev/null
+++ b/tests/lang/test_sycl.py
@@ -0,0 +1,41 @@
+import pytest
+
+
+from pystencilssfg.lang.cpp.sycl_accessor import sycl_accessor_ref
+import pystencils as ps
+
+
+@pytest.mark.parametrize("data_type", ["double", "float"])
+@pytest.mark.parametrize("dim", [1, 2, 3])
+def test_spatial_field(data_type, dim):
+ f = ps.fields(f"f:{data_type}[{dim}D]")
+ ref = sycl_accessor_ref(f)
+ assert (
+ f"sycl::accessor< {data_type}, {dim} > &" in ref.as_variable().name_and_type()
+ )
+
+
+@pytest.mark.parametrize("data_type", ["double", "float"])
+def test_to_large_dim(data_type):
+ dim = 4
+ f = ps.fields(f"f:{data_type}[{dim}D]")
+ with pytest.raises(ValueError):
+ sycl_accessor_ref(f)
+
+
+@pytest.mark.parametrize("data_type", ["double", "float"])
+@pytest.mark.parametrize("spatial_dim", [1, 2, 3])
+@pytest.mark.parametrize("index_dims", [1, 2, 3])
+def test_index_field(data_type, spatial_dim, index_dims):
+ index_shape = ("19",) * index_dims
+ total_dims = spatial_dim + index_dims
+ f = ps.fields(f"f({', '.join(index_shape)}):{data_type}[{spatial_dim}D]")
+ if total_dims <= 3:
+ ref = sycl_accessor_ref(f)
+ assert (
+ f"sycl::accessor< {data_type}, {total_dims} > &"
+ in ref.as_variable().name_and_type()
+ )
+ else:
+ with pytest.raises(ValueError):
+ sycl_accessor_ref(f)
--
GitLab
From 77316bdfcba3e152e7c629fcd3defd0a3ea3f31d Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Thu, 28 Nov 2024 13:54:14 +0100
Subject: [PATCH 10/13] added a test for the parallel_for
---
tests/extensions/test_sycl.py | 64 +++++++++++++++++++
.../{test_sycl.py => test_sycl_accessor.py} | 0
2 files changed, 64 insertions(+)
create mode 100644 tests/extensions/test_sycl.py
rename tests/lang/{test_sycl.py => test_sycl_accessor.py} (100%)
diff --git a/tests/extensions/test_sycl.py b/tests/extensions/test_sycl.py
new file mode 100644
index 0000000..db99278
--- /dev/null
+++ b/tests/extensions/test_sycl.py
@@ -0,0 +1,64 @@
+import pytest
+from pystencilssfg import SourceFileGenerator
+import pystencilssfg.extensions.sycl as sycl
+import pystencils as ps
+from pystencilssfg import SfgContext
+
+
+def test_parallel_for_1_kernels():
+ sfg = sycl.SyclComposer(SfgContext())
+ data_type = "double"
+ dim = 2
+ f, g, h, i = ps.fields(f"f,g,h,i:{data_type}[{dim}D]")
+ assignement_1 = ps.Assignment(f.center(), g.center())
+ assignement_2 = ps.Assignment(h.center(), i.center())
+
+ config = ps.CreateKernelConfig(target=ps.Target.SYCL)
+ kernel_1 = sfg.kernels.create(assignement_1, "kernel_1", config)
+ kernel_2 = sfg.kernels.create(assignement_2, "kernel_2", config)
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(dim, "range")
+ cgh.parallel_for(rang)(
+ sfg.call(kernel_1),
+ sfg.call(kernel_2),
+ )
+
+
+def test_parallel_for_2_kernels():
+ sfg = sycl.SyclComposer(SfgContext())
+ data_type = "double"
+ dim = 2
+ f, g, h, i = ps.fields(f"f,g,h,i:{data_type}[{dim}D]")
+ assignement_1 = ps.Assignment(f.center(), g.center())
+ assignement_2 = ps.Assignment(h.center(), i.center())
+
+ config = ps.CreateKernelConfig(target=ps.Target.SYCL)
+ kernel_1 = sfg.kernels.create(assignement_1, "kernel_1", config)
+ kernel_2 = sfg.kernels.create(assignement_2, "kernel_2", config)
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(dim, "range")
+ cgh.parallel_for(rang)(
+ sfg.call(kernel_1),
+ sfg.call(kernel_2),
+ )
+
+
+def test_parallel_for_2_kernels_fail():
+ sfg = sycl.SyclComposer(SfgContext())
+ data_type = "double"
+ dim = 2
+ f, g = ps.fields(f"f,g:{data_type}[{dim}D]")
+ h, i = ps.fields(f"h,i:{data_type}[{dim-1}D]")
+ assignement_1 = ps.Assignment(f.center(), g.center())
+ assignement_2 = ps.Assignment(h.center(), i.center())
+
+ config = ps.CreateKernelConfig(target=ps.Target.SYCL)
+ kernel_1 = sfg.kernels.create(assignement_1, "kernel_1", config)
+ kernel_2 = sfg.kernels.create(assignement_2, "kernel_2", config)
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(dim, "range")
+ with pytest.raises(ValueError):
+ cgh.parallel_for(rang)(
+ sfg.call(kernel_1),
+ sfg.call(kernel_2),
+ )
diff --git a/tests/lang/test_sycl.py b/tests/lang/test_sycl_accessor.py
similarity index 100%
rename from tests/lang/test_sycl.py
rename to tests/lang/test_sycl_accessor.py
--
GitLab
From 0ee586eff5cecfc74c996446514a50f355296dfc Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Tue, 3 Dec 2024 14:48:22 +0100
Subject: [PATCH 11/13] updated the integration test to new config api
---
integration/test_sycl_buffer.py | 7 +++----
1 file changed, 3 insertions(+), 4 deletions(-)
diff --git a/integration/test_sycl_buffer.py b/integration/test_sycl_buffer.py
index 7839b9f..a1e52eb 100644
--- a/integration/test_sycl_buffer.py
+++ b/integration/test_sycl_buffer.py
@@ -1,15 +1,14 @@
from pystencils import Target, CreateKernelConfig, no_jit
from lbmpy import create_lb_update_rule, LBMOptimisation
-from pystencilssfg import SourceFileGenerator, SfgConfiguration, SfgOutputMode
+from pystencilssfg import SourceFileGenerator, SfgConfig, OutputMode
from pystencilssfg.lang.cpp.sycl_accessor import sycl_accessor_ref
import pystencilssfg.extensions.sycl as sycl
from itertools import chain
-sfg_config = SfgConfiguration(
+sfg_config = SfgConfig(
output_directory="out/test_sycl_buffer",
outer_namespace="gen_code",
- impl_extension="ipp",
- output_mode=SfgOutputMode.INLINE,
+ output_mode=OutputMode.INLINE,
)
with SourceFileGenerator(sfg_config) as sfg:
--
GitLab
From f8ed8ad8d810be016badde83279265a61e187f3f Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Thu, 5 Dec 2024 08:45:37 +0100
Subject: [PATCH 12/13] Added a generator script to test the sycl accessors
---
.../scripts/TestSyclBuffer.py | 35 +++++++++++++++++++
.../test_generator_scripts.py | 18 +++++++---
2 files changed, 49 insertions(+), 4 deletions(-)
create mode 100644 tests/generator_scripts/scripts/TestSyclBuffer.py
diff --git a/tests/generator_scripts/scripts/TestSyclBuffer.py b/tests/generator_scripts/scripts/TestSyclBuffer.py
new file mode 100644
index 0000000..d041756
--- /dev/null
+++ b/tests/generator_scripts/scripts/TestSyclBuffer.py
@@ -0,0 +1,35 @@
+import pystencils as ps
+import sympy as sp
+from pystencilssfg import SourceFileGenerator
+from pystencilssfg.lang.cpp.sycl_accessor import sycl_accessor_ref
+import pystencilssfg.extensions.sycl as sycl
+
+
+with SourceFileGenerator() as sfg:
+ sfg = sycl.SyclComposer(sfg)
+
+ u_src, u_dst, f = ps.fields("u_src, u_dst, f : double[2D]", layout="fzyx")
+ h = sp.Symbol("h")
+
+ jacobi_update = [
+ ps.Assignment(
+ u_dst.center(),
+ (h**2 * f[0, 0] + u_src[1, 0] + u_src[-1, 0] + u_src[0, 1] + u_src[0, -1])
+ / 4,
+ )
+ ]
+
+ kernel_config = ps.CreateKernelConfig(target=ps.Target.SYCL)
+ jacobi_kernel = sfg.kernels.create(jacobi_update, config=kernel_config)
+
+ cgh = sfg.sycl_handler("handler")
+ rang = sfg.sycl_range(2, "range")
+ mappings = [
+ sfg.map_field(u_src, sycl_accessor_ref(u_src)),
+ sfg.map_field(u_dst, sycl_accessor_ref(u_dst)),
+ sfg.map_field(f, sycl_accessor_ref(f)),
+ ]
+
+ sfg.function("jacobiUpdate")(
+ cgh.parallel_for(rang)(*mappings, sfg.call(jacobi_kernel)),
+ )
diff --git a/tests/generator_scripts/test_generator_scripts.py b/tests/generator_scripts/test_generator_scripts.py
index 16adba2..316e0e0 100644
--- a/tests/generator_scripts/test_generator_scripts.py
+++ b/tests/generator_scripts/test_generator_scripts.py
@@ -74,7 +74,7 @@ SCRIPTS = [
"--sfg-file-extensionss",
".c++,.h++",
),
- should_fail=True
+ should_fail=True,
),
ScriptInfo.make(
"TestExtraCommandLineArgs",
@@ -85,13 +85,19 @@ SCRIPTS = [
"--precision",
"float32",
"test1",
- "test2"
+ "test2",
),
),
ScriptInfo.make("Structural", ("hpp", "cpp")),
ScriptInfo.make("SimpleJacobi", ("hpp", "cpp"), compilable_output="cpp"),
ScriptInfo.make("SimpleClasses", ("hpp", "cpp")),
ScriptInfo.make("Variables", ("hpp", "cpp"), compilable_output="cpp"),
+ ScriptInfo.make(
+ "TestSyclBuffer",
+ ("hpp", "cpp"),
+ compilable_output="cpp" if shutil.which("icpx") else None,
+ compile_cmd="icpx -fsycl -std=c++20" if shutil.which("icpx") else "",
+ ),
]
@@ -113,13 +119,17 @@ def test_generator_script(script_info: ScriptInfo):
shutil.rmtree(output_dir)
os.makedirs(output_dir, exist_ok=True)
- args = ["python", script_file, "--sfg-output-dir", output_dir] + list(script_info.args)
+ args = ["python", script_file, "--sfg-output-dir", output_dir] + list(
+ script_info.args
+ )
result = subprocess.run(args)
if script_info.should_fail:
if result.returncode == 0:
- pytest.fail(f"Generator script {script_name} was supposed to fail, but didn't.")
+ pytest.fail(
+ f"Generator script {script_name} was supposed to fail, but didn't."
+ )
return
if result.returncode != 0:
--
GitLab
From f02641fe3be5ccff79dcf01b4611aec11c06e40f Mon Sep 17 00:00:00 2001
From: Christoph Alt <christoph.alt@fau.de>
Date: Thu, 5 Dec 2024 08:50:03 +0100
Subject: [PATCH 13/13] using get_dtype() to get the full typename
---
tests/lang/test_sycl_accessor.py | 11 +++--------
1 file changed, 3 insertions(+), 8 deletions(-)
diff --git a/tests/lang/test_sycl_accessor.py b/tests/lang/test_sycl_accessor.py
index 88f958f..d2af295 100644
--- a/tests/lang/test_sycl_accessor.py
+++ b/tests/lang/test_sycl_accessor.py
@@ -10,13 +10,11 @@ import pystencils as ps
def test_spatial_field(data_type, dim):
f = ps.fields(f"f:{data_type}[{dim}D]")
ref = sycl_accessor_ref(f)
- assert (
- f"sycl::accessor< {data_type}, {dim} > &" in ref.as_variable().name_and_type()
- )
+ assert f"sycl::accessor< {data_type}, {dim} > &" in str(ref.get_dtype())
@pytest.mark.parametrize("data_type", ["double", "float"])
-def test_to_large_dim(data_type):
+def test_too_large_dim(data_type):
dim = 4
f = ps.fields(f"f:{data_type}[{dim}D]")
with pytest.raises(ValueError):
@@ -32,10 +30,7 @@ def test_index_field(data_type, spatial_dim, index_dims):
f = ps.fields(f"f({', '.join(index_shape)}):{data_type}[{spatial_dim}D]")
if total_dims <= 3:
ref = sycl_accessor_ref(f)
- assert (
- f"sycl::accessor< {data_type}, {total_dims} > &"
- in ref.as_variable().name_and_type()
- )
+ assert f"sycl::accessor< {data_type}, {total_dims} > &" in str(ref.get_dtype())
else:
with pytest.raises(ValueError):
sycl_accessor_ref(f)
--
GitLab