Skip to content
Snippets Groups Projects
Commit ef144a22 authored by Christoph Alt's avatar Christoph Alt
Browse files

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
parent 9672cf57
No related branches found
No related tags found
1 merge request!3Add support for sycl accessors
import math
from ...lang import SrcField, IFieldExtraction from ...lang import SrcField, IFieldExtraction
from ...ir.source_components import SfgHeaderInclude from ...ir.source_components import SfgHeaderInclude
from typing import Sequence
from pystencils import Field from pystencils import Field
from pystencils.types import ( from pystencils.types import (
...@@ -17,22 +15,17 @@ class SyclAccessor(SrcField): ...@@ -17,22 +15,17 @@ class SyclAccessor(SrcField):
self, self,
T: PsType, T: PsType,
dimensions: int, dimensions: int,
index_shape: Sequence[int],
reference: bool = False, reference: bool = False,
): ):
cpp_typestr = T.c_string() 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") raise ValueError("sycl accessors can only have dims 1, 2 or 3")
typestring = ( typestring = (
f"sycl::accessor< {cpp_typestr}, {dimensions} > {'&' if reference else ''}" f"sycl::accessor< {cpp_typestr}, {dimensions} > {'&' if reference else ''}"
) )
super().__init__(PsCustomType(typestring)) super().__init__(PsCustomType(typestring))
self._dim = dimensions
self._spatial_dimensions = dimensions self._inner_stride = 1
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 @property
def required_includes(self) -> set[SfgHeaderInclude]: def required_includes(self) -> set[SfgHeaderInclude]:
...@@ -49,7 +42,7 @@ class SyclAccessor(SrcField): ...@@ -49,7 +42,7 @@ class SyclAccessor(SrcField):
) )
def size(self, coordinate: int) -> AugExpr | None: def size(self, coordinate: int) -> AugExpr | None:
if coordinate > accessor._spatial_dimensions: if coordinate > accessor._dim:
return None return None
else: else:
return AugExpr.format( return AugExpr.format(
...@@ -57,36 +50,30 @@ class SyclAccessor(SrcField): ...@@ -57,36 +50,30 @@ class SyclAccessor(SrcField):
) )
def stride(self, coordinate: int) -> AugExpr | None: def stride(self, coordinate: int) -> AugExpr | None:
if coordinate > accessor._total_dimensions_: if coordinate > accessor._dim:
return None return None
elif coordinate >= accessor._spatial_dimensions - 1: elif coordinate == accessor._dim - 1:
start = (coordinate - accessor._spatial_dimensions) + 1 return AugExpr.format("{}", accessor._inner_stride)
return AugExpr.format(
"{}", math.prod(accessor._index_shape[start:])
)
else: else:
exprs = [] exprs = []
args = [] args = []
for d in range(coordinate + 1, accessor._spatial_dimensions): for d in range(coordinate + 1, accessor._dim):
args.extend([accessor, d]) args.extend([accessor, d])
exprs.append("{}.get_range().get({})") exprs.append("{}.get_range().get({})")
expr = " * ".join(exprs) expr = " * ".join(exprs)
expr += " * {}" expr += " * {}"
return AugExpr.format(expr, *args, accessor._index_size) return AugExpr.format(expr, *args, accessor._inner_stride)
return Extraction() return Extraction()
def sycl_accessor_ref(field: Field): def sycl_accessor_ref(field: Field):
"""Creates a `sycl::accessor &` for a given pystencils field.""" """Creates a `sycl::accessor &` for a given pystencils field."""
# Sycl accesors allow only at max 3 dimensions: # Sycl Accessor do not expose information about strides, so the linearization is like here
# So also mapping the index dimens to the sycl accesor we only can have 2D LBM stuff # https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_multi_dimensional_objects_and_linearization
# 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( return SyclAccessor(
field.dtype, field.dtype,
field.spatial_dimensions, field.spatial_dimensions + field.index_dimensions,
field.index_shape,
reference=True, reference=True,
).var(field.name) ).var(field.name)
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)
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment