Skip to content
Snippets Groups Projects
Commit 422a2494 authored by Frederik Hennig's avatar Frederik Hennig
Browse files

add dry-dock tests for SYCL codegen

parent 1dff4610
No related branches found
No related tags found
1 merge request!453Clean up kernel creation code in `codegen.driver`. Fix and dry-test SYCL codegen.
Pipeline #74264 passed
......@@ -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"
......
"""
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
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment