diff --git a/src/pystencils/codegen/driver.py b/src/pystencils/codegen/driver.py index eaac393ef61bf937835b0b3a5043aa4395edde4d..b8f9c71015765e638ddca22278dfa15c0e5bcaa1 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 0000000000000000000000000000000000000000..b6907c9965b4b80a5b97865063170dfdc3654615 --- /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