Skip to content

Commit

Permalink
Set extra compiler options inside KernelCompiler pipeline.
Browse files Browse the repository at this point in the history
    - We should not generate a cpython or a cfunc wrapper for a
      kernel decorated function.
    - The no_compile is also set to True as the kernel LLVM IR
      will not be compiled to an executable object code by Numba.
      Instead, numba-dpex will compile the LLVM IR to SPIR-V and
      generate a sycl::kernel_bundle.
  • Loading branch information
Diptorup Deb committed Oct 12, 2023
1 parent d254ee4 commit 9576aa5
Showing 1 changed file with 11 additions and 2 deletions.
13 changes: 11 additions & 2 deletions numba_dpex/core/pipelines/kernel_compiler.py
Original file line number Diff line number Diff line change
Expand Up @@ -140,7 +140,6 @@ def define_nopython_lowering_pipeline(state, name="dpex_kernel_lowering"):
)
pm.add_pass(IRLegalization, "ensure IR is legal prior to lowering")

# lower
# NativeLowering has some issue with freevar ambiguity,
# therefore, we are using QualNameDisambiguationLowering instead
# numba-dpex github issue: https://github.com/IntelPython/numba-dpex/issues/898
Expand Down Expand Up @@ -173,12 +172,22 @@ def define_nopython_pipeline(state, name="dpex_kernel_nopython"):


class KernelCompiler(CompilerBase):
"""Dpex's kernel compilation pipeline."""
"""Dpex's kernel compilation pipeline."""

def define_pipelines(self):
pms = []
if not self.state.flags.force_pyobject:
pms.append(_KernelPassBuilder.define_nopython_pipeline(self.state))
if self.state.status.can_fallback or self.state.flags.force_pyobject:
raise UnsupportedCompilationModeError()

# Compile the kernel without generating a cpython or a cfunc wrapper
self.state.flags.no_cpython_wrapper = True
self.state.flags.no_cfunc_wrapper = True
# The pass pipeline does not generate an executable when compiling a
# kernel function. Instead, the
# kernel_dispatcher._KernelCompiler.compile generates the executable in
# the form of a host callable launcher function
self.state.flags.no_compile = True

return pms

0 comments on commit 9576aa5

Please sign in to comment.