diff --git a/src/pystencilssfg/composer/basic_composer.py b/src/pystencilssfg/composer/basic_composer.py index 49b6c73e8bd06b0d5fb44402b8af285362072b29..ff1ab2df679aeefd2112d539b90dd59cf93e015b 100644 --- a/src/pystencilssfg/composer/basic_composer.py +++ b/src/pystencilssfg/composer/basic_composer.py @@ -88,11 +88,16 @@ SequencerArg: TypeAlias = tuple | ExprLike | SfgCallTreeNode | SfgNodeBuilder class KernelsAdder: """Handle on a kernel namespace that permits registering kernels.""" - def __init__(self, ctx: SfgContext, loc: SfgNamespaceBlock): - self._ctx = ctx - self._loc = loc - assert isinstance(loc.namespace, SfgKernelNamespace) - self._kernel_namespace = loc.namespace + def __init__(self, cursor: SfgCursor, knamespace: SfgKernelNamespace): + self._cursor = cursor + self._kernel_namespace = knamespace + self._inline: bool = False + self._loc: SfgNamespaceBlock | None = None + + def inline(self) -> KernelsAdder: + """Generate kernel definitions ``inline`` in the header file.""" + self._inline = True + return self def add(self, kernel: Kernel, name: str | None = None): """Adds an existing pystencils AST to this namespace. @@ -111,14 +116,22 @@ class KernelsAdder: if name is not None: kernel.name = kernel_name - khandle = SfgKernelHandle(kernel_name, self._kernel_namespace, kernel) + khandle = SfgKernelHandle( + kernel_name, self._kernel_namespace, kernel, inline=self._inline + ) self._kernel_namespace.add_kernel(khandle) - self._loc.elements.append(SfgEntityDef(khandle)) + loc = self._get_loc() + loc.elements.append(SfgEntityDef(khandle)) for header in kernel.required_headers: - assert self._ctx.impl_file is not None - self._ctx.impl_file.includes.append(HeaderFile.parse(header)) + hfile = HeaderFile.parse(header) + if self._inline: + self._cursor.context.header_file.includes.append(hfile) + else: + impl_file = self._cursor.context.impl_file + assert impl_file is not None + impl_file.includes.append(hfile) return khandle @@ -147,6 +160,18 @@ class KernelsAdder: kernel = create_kernel(assignments, config=config) return self.add(kernel) + def _get_loc(self) -> SfgNamespaceBlock: + if self._loc is None: + kns_block = SfgNamespaceBlock(self._kernel_namespace) + + if self._inline: + self._cursor.write_header(kns_block) + else: + self._cursor.write_impl(kns_block) + + self._loc = kns_block + return self._loc + class SfgBasicComposer(SfgIComposer): """Composer for basic source components, and base class for all composer mix-ins.""" @@ -281,9 +306,10 @@ class SfgBasicComposer(SfgIComposer): f"The existing entity {kns.fqname} is not a kernel namespace" ) - kns_block = SfgNamespaceBlock(kns) - self._cursor.write_impl(kns_block) - return KernelsAdder(self._ctx, kns_block) + kadder = KernelsAdder(self._cursor, kns) + if self._ctx.impl_file is None: + kadder.inline() + return kadder def include(self, header: str | HeaderFile, private: bool = False): """Include a header file. @@ -356,6 +382,9 @@ class SfgBasicComposer(SfgIComposer): ) seq.returns(return_type) + if self._ctx.impl_file is None: + seq.inline() + return seq def call(self, kernel_handle: SfgKernelHandle) -> SfgCallTreeNode: diff --git a/src/pystencilssfg/composer/class_composer.py b/src/pystencilssfg/composer/class_composer.py index 7787150c5bffd4e7332aecd2ff7a4a66cac0adc6..5dbb4c65d53c9305225239cdaf9753413d07d106 100644 --- a/src/pystencilssfg/composer/class_composer.py +++ b/src/pystencilssfg/composer/class_composer.py @@ -253,7 +253,10 @@ class SfgClassComposer(SfgComposerMixIn): name: The method name """ - return SfgMethodSequencer(self._cursor, name) + seq = SfgMethodSequencer(self._cursor, name) + if self._ctx.impl_file is None: + seq.inline() + return seq # INTERNALS diff --git a/src/pystencilssfg/context.py b/src/pystencilssfg/context.py index 199c678ba28e449f42b29c56147b9a4fd0d523bb..1622a1e3bd33259cd89eef171ad043c4ea9ef536 100644 --- a/src/pystencilssfg/context.py +++ b/src/pystencilssfg/context.py @@ -115,6 +115,10 @@ class SfgCursor: else: self._loc[f] = f.elements + @property + def context(self) -> SfgContext: + return self._ctx + @property def current_namespace(self) -> SfgNamespace: return self._cur_namespace diff --git a/src/pystencilssfg/emission/file_printer.py b/src/pystencilssfg/emission/file_printer.py index 765bf70550504fd499746504236f1adaa224664a..648e41971336fac709f21848f729b025de27ff36 100644 --- a/src/pystencilssfg/emission/file_printer.py +++ b/src/pystencilssfg/emission/file_printer.py @@ -27,9 +27,7 @@ from ..config import CodeStyle class SfgFilePrinter: def __init__(self, code_style: CodeStyle) -> None: self._code_style = code_style - self._kernel_printer = CAstPrinter( - indent_width=code_style.get_option("indent_width") - ) + self._indent_width = code_style.get_option("indent_width") def __call__(self, file: SfgSourceFile) -> str: code = "" @@ -86,7 +84,11 @@ class SfgFilePrinter: ) -> str: match declared_entity: case SfgKernelHandle(kernel): - return self._kernel_printer.print_signature(kernel) + ";" + kernel_printer = CAstPrinter( + indent_width=self._indent_width, + func_prefix="inline" if declared_entity.inline else None, + ) + return kernel_printer.print_signature(kernel) + ";" case SfgFunction(name, _, params) | SfgMethod(name, _, params): return self._func_signature(declared_entity, inclass) + ";" @@ -113,7 +115,11 @@ class SfgFilePrinter: ) -> str: match defined_entity: case SfgKernelHandle(kernel): - return self._kernel_printer(kernel) + kernel_printer = CAstPrinter( + indent_width=self._indent_width, + func_prefix="inline" if defined_entity.inline else None, + ) + return kernel_printer(kernel) case SfgFunction(name, tree, params) | SfgMethod(name, tree, params): sig = self._func_signature(defined_entity, inclass) diff --git a/src/pystencilssfg/ir/entities.py b/src/pystencilssfg/ir/entities.py index 40abb148d90eb35300afd0af4e25d2f8df1c091e..0edde2209a7ec7bab102de33202aecd3011bfd8a 100644 --- a/src/pystencilssfg/ir/entities.py +++ b/src/pystencilssfg/ir/entities.py @@ -141,12 +141,20 @@ class SfgKernelHandle(SfgCodeEntity): __match_args__ = ("kernel", "parameters") - def __init__(self, name: str, namespace: SfgKernelNamespace, kernel: Kernel): + def __init__( + self, + name: str, + namespace: SfgKernelNamespace, + kernel: Kernel, + inline: bool = False, + ): super().__init__(name, namespace) self._kernel = kernel self._parameters = [SfgKernelParamVar(p) for p in kernel.parameters] + self._inline: bool = inline + self._scalar_params: set[SfgVar] = set() self._fields: set[Field] = set() @@ -176,6 +184,10 @@ class SfgKernelHandle(SfgCodeEntity): """Underlying pystencils kernel object""" return self._kernel + @property + def inline(self) -> bool: + return self._inline + class SfgKernelNamespace(SfgNamespace): """A namespace grouping together a number of kernels.""" diff --git a/tests/generator_scripts/index.yaml b/tests/generator_scripts/index.yaml index 5e2db9a65b6ccc13f555c160144be93e652a04c1..e7db3479e35fc22ac204e5efbbbca7afefe8b9ad 100644 --- a/tests/generator_scripts/index.yaml +++ b/tests/generator_scripts/index.yaml @@ -53,6 +53,16 @@ ComposerFeatures: - regex: >- \[\[nodiscard\]\]\s*static\s*double\s*geometric\(\s*double\s*q,\s*uint64_t\s*k\) +ComposerHeaderOnly: + sfg-args: + header-only: true + expect-code: + hpp: + - regex: >- + inline\sint32_t\stwice\( + - regex: >- + inline\svoid\skernel\( + Conditionals: expect-code: cpp: @@ -84,10 +94,11 @@ MdSpanLbStreaming: SyclKernels: sfg-args: - output-mode: inline - file-extensions: [hpp, ipp] + header-only: true expect-code: - ipp: + hpp: + - regex: >- + inline\svoid\skernel\( - regex: >- cgh\.parallel_for\(range,\s*\[=\]\s*\(const\s+sycl::item<\s*2\s*>\s+sycl_item\s*\)\s*\{\s*kernels::kernel\(.*\);\s*\}\); diff --git a/tests/generator_scripts/source/ComposerHeaderOnly.harness.cpp b/tests/generator_scripts/source/ComposerHeaderOnly.harness.cpp new file mode 100644 index 0000000000000000000000000000000000000000..f6c549a3a33d434fd0bb8865140f1edf29ed9983 --- /dev/null +++ b/tests/generator_scripts/source/ComposerHeaderOnly.harness.cpp @@ -0,0 +1,18 @@ +#include "ComposerHeaderOnly.hpp" + +#include <vector> + +#undef NDEBUG +#include <cassert> + +int main(void) { + assert( twice(13) == 26 ); + + std::vector< int64_t > arr { 1, 2, 3, 4, 5, 6 }; + twiceKernel(arr); + + std::vector< int64_t > expected { 2, 4, 6, 8, 10, 12 }; + assert ( arr == expected ); + + return 0; +} diff --git a/tests/generator_scripts/source/ComposerHeaderOnly.py b/tests/generator_scripts/source/ComposerHeaderOnly.py new file mode 100644 index 0000000000000000000000000000000000000000..0bd1a07879855b9a2af75e493fa84e92632a70ff --- /dev/null +++ b/tests/generator_scripts/source/ComposerHeaderOnly.py @@ -0,0 +1,26 @@ +from pystencilssfg import SourceFileGenerator, SfgConfig +from pystencilssfg.lang.cpp import std +import pystencils as ps + +cfg = SfgConfig(header_only=True) + +with SourceFileGenerator(cfg) as sfg: + n = sfg.var("n", "int32") + + # Should be automatically marked inline + sfg.function("twice").returns("int32")( + sfg.expr("return 2 * {};", n) + ) + + # Inline kernel + + arr = ps.fields("arr: int64[1D]") + vec = std.vector.from_field(arr) + + asm = ps.Assignment(arr(0), 2 * arr(0)) + khandle = sfg.kernels.create(asm) + + sfg.function("twiceKernel")( + sfg.map_field(arr, vec), + sfg.call(khandle) + )