From 422a2494c1ae20f1d28f8cd8e2a52faf67e75c1f Mon Sep 17 00:00:00 2001
From: Frederik Hennig <frederik.hennig@fau.de>
Date: Wed, 19 Feb 2025 10:27:40 +0100
Subject: [PATCH] add dry-dock tests for SYCL codegen
---
src/pystencils/codegen/driver.py | 25 ++++++-------
tests/kernelcreation/test_sycl_codegen.py | 45 +++++++++++++++++++++++
2 files changed, 57 insertions(+), 13 deletions(-)
create mode 100644 tests/kernelcreation/test_sycl_codegen.py
diff --git a/src/pystencils/codegen/driver.py b/src/pystencils/codegen/driver.py
index eaac393ef..b8f9c7101 100644
--- a/src/pystencils/codegen/driver.py
+++ b/src/pystencils/codegen/driver.py
@@ -436,23 +436,11 @@ class DefaultKernelCreationDriver:
f"No platform is currently available for CPU target {self._target}"
)
- elif Target._GPU in self._target:
+ elif self._target.is_gpu():
gpu_opts = self._cfg.gpu
omit_range_check: bool = gpu_opts.get_option("omit_range_check")
match self._target:
- case Target.SYCL:
- from ..backend.platforms import SyclPlatform
-
- auto_block_size: bool = self._cfg.sycl.get_option(
- "automatic_block_size"
- )
-
- return SyclPlatform(
- self._ctx,
- omit_range_check=omit_range_check,
- automatic_block_size=auto_block_size,
- )
case Target.CUDA:
from ..backend.platforms import CudaPlatform
@@ -467,6 +455,17 @@ class DefaultKernelCreationDriver:
omit_range_check=omit_range_check,
thread_mapping=thread_mapping,
)
+ elif self._target == Target.SYCL:
+ from ..backend.platforms import SyclPlatform
+
+ auto_block_size: bool = self._cfg.sycl.get_option("automatic_block_size")
+ omit_range_check = self._cfg.gpu.get_option("omit_range_check")
+
+ return SyclPlatform(
+ self._ctx,
+ omit_range_check=omit_range_check,
+ automatic_block_size=auto_block_size,
+ )
raise NotImplementedError(
f"Code generation for target {self._target} not implemented"
diff --git a/tests/kernelcreation/test_sycl_codegen.py b/tests/kernelcreation/test_sycl_codegen.py
new file mode 100644
index 000000000..b6907c996
--- /dev/null
+++ b/tests/kernelcreation/test_sycl_codegen.py
@@ -0,0 +1,45 @@
+"""
+Since we don't have a JIT compiler for SYCL, these tests can only
+perform dry-dock testing.
+If the SYCL target should ever become non-experimental, we need to
+find a way to properly test SYCL kernels in execution.
+
+These tests primarily check that the code generation driver runs
+successfully for the SYCL target.
+"""
+
+import sympy as sp
+from pystencils import (
+ create_kernel,
+ Target,
+ fields,
+ Assignment,
+ CreateKernelConfig,
+)
+
+
+def test_sycl_kernel_static():
+ src, dst = fields("src, dst: [2D]")
+ asm = Assignment(dst.center(), sp.sin(src.center()) + sp.cos(src.center()))
+
+ cfg = CreateKernelConfig(target=Target.SYCL)
+ kernel = create_kernel(asm, cfg)
+
+ code_string = kernel.get_c_code()
+
+ assert "sycl::id< 2 >" in code_string
+ assert "sycl::sin(" in code_string
+ assert "sycl::cos(" in code_string
+
+
+def test_sycl_kernel_manual_block_size():
+ src, dst = fields("src, dst: [2D]")
+ asm = Assignment(dst.center(), sp.sin(src.center()) + sp.cos(src.center()))
+
+ cfg = CreateKernelConfig(target=Target.SYCL)
+ cfg.sycl.automatic_block_size = False
+ kernel = create_kernel(asm, cfg)
+
+ code_string = kernel.get_c_code()
+
+ assert "sycl::nd_item< 2 >" in code_string
--
GitLab