Exceptions

The helion.exc module provides a exception hierarchy for error handling and diagnostics.

Overview

Helion’s exception system provides detailed error messages with automatic source location tracking. All exceptions inherit from Base and provide formatted error reports that help identify exactly where and why compilation failed.

Exception Hierarchy

Base Classes

class helion.exc.Base[source]

Bases: RuntimeError

report()[source]
Return type:

str

class helion.exc.BaseError(*args, **kwargs)[source]

Bases: _FixedMessage

Parameters:
message = 'An error occurred.'
report()[source]
Return type:

str

class helion.exc.BaseWarning(*args, **kwargs)[source]

Bases: _FixedMessage

Parameters:
message = 'A warning occurred.'
report()[source]
Return type:

str

Kernel Context Errors

These exceptions occur when Helion language functions are used incorrectly with respect to kernel context:

class helion.exc.NotInsideKernel(*args, **kwargs)[source]
Parameters:

Raised when helion.language.* functions are called outside of a kernel context.

Example:

import helion.language as hl

# This will raise NotInsideKernel
result = hl.zeros([10])  # Called outside @helion.kernel
message = 'Functions found in helion.language.* must be called from inside a kernel. Did you forget the @helion.kernel decorator?'
class helion.exc.NotAllowedOnDevice(*args, **kwargs)[source]
Parameters:

Raised when host-only operations are used inside hl.tile() or hl.grid() loops.

message = 'The statement {} is not allowed inside the `hl.tile` or `hl.grid` loop.'
class helion.exc.DeviceAPIOnHost(*args, **kwargs)[source]
Parameters:

Raised when device-only APIs are called in host context.

message = '{} is only allowed inside the `hl.tile` or `hl.grid` loop.'
class helion.exc.CantReadOnDevice(*args, **kwargs)[source]
Parameters:

Raised when attempting to read host variables from device code.

message = 'Cannot read {0!s} inside the `hl.tile` or `hl.grid` loop.'
class helion.exc.NotAllowedInHelperFunction(*args, **kwargs)[source]
Parameters:

Raised when operations requiring kernel context are used in helper functions.

message = 'This operation is not allowed inside helper functions. It requires kernel context.'

Loop and Control Flow Errors

class helion.exc.LoopFunctionNotInFor(*args, **kwargs)[source]
Parameters:

Raised when hl.tile() or hl.grid() are called outside for loops.

Correct usage:

for i in hl.grid(size):  # Correct
    pass

i = hl.grid(size)        # Raises LoopFunctionNotInFor
message = '{0} must be called from a for loop, e.g. `for ... in {0}(...):'
class helion.exc.InvalidDeviceForLoop(*args, **kwargs)[source]
Parameters:

Raised for invalid for loop constructs on device (must use hl.tile/hl.grid).

message = 'For loops on device must use `hl.tile` or `hl.grid`, got {0!s}.'
class helion.exc.NestedDeviceLoopsConflict(*args, **kwargs)[source]
Parameters:

Raised when nested device loops have conflicting block sizes.

message = 'Nested device loops must have distinct block sizes.'
class helion.exc.DeviceLoopElseBlock(*args, **kwargs)[source]
Parameters:

Raised when for…else blocks are used in device loops.

message = 'for...else block is not allowed in a {0} device loop.'
class helion.exc.NestedGridLoop(*args, **kwargs)[source]
Parameters:

Raised when grid loops are not at function top level.

message = 'Grid loops must be at the top level of a function.'
class helion.exc.TopLevelStatementBetweenLoops(*args, **kwargs)[source]
Parameters:

Raised when statements appear between top-level loops.

message = 'Statements cannot appear between top level loops.'
class helion.exc.LoopDependencyError(*args, **kwargs)[source]
Parameters:

Raised when writing to variables across loop iterations creates dependencies.

message = "Loop dependency detected: '{0}' was written in a previous loop."

Tile and Indexing Errors

class helion.exc.IncorrectTileUsage(*args, **kwargs)[source]
Parameters:

Raised when tiles are used outside tensor indexing or hl.* operations.

message = 'Tiles can only be used in tensor indexing (`x[tile]`) or in `hl.*` ops (e.g. `hl.zeros(tile)`), used in {}'
class helion.exc.FailedToUnpackTile(*args, **kwargs)[source]
Parameters:

Raised when tuple unpacking fails for single tile.

message = 'Failed to unpack a tile into a tuple assignment. Expected a sequence, but got a single tile. Did you mix up `hl.tile(x)` and `hl.tile([x])`?'
class helion.exc.OverpackedTile(*args, **kwargs)[source]
Parameters:

Raised when tile is wrapped in container when indexing.

message = 'Got a tile wrapped inside a container when indexing a tensor: {0!s}\nDid you mix up `hl.tile([x])` and `hl.tile(x)`?'
class helion.exc.RankMismatch(expected_ndim, actual_ndim, shape_info='')[source]
Parameters:
  • expected_ndim (int)

  • actual_ndim (int)

  • shape_info (str)

Raised when tensor rank doesn’t match indexing dimensions.

Example:

x = torch.randn(10, 20)  # 2D tensor
for i in hl.grid(10):
    y = x[i, j, k]       # Raises RankMismatch - too many indices
message = 'Expected ndim={expected_ndim}, but got ndim={actual_ndim}{shape_part}. You have {direction}.'
__init__(expected_ndim, actual_ndim, shape_info='')[source]
Parameters:
  • expected_ndim (int)

  • actual_ndim (int)

  • shape_info (str)

class helion.exc.InvalidIndexingType(*args, **kwargs)[source]
Parameters:

Raised for invalid types in tensor subscripts.

message = 'Expected tile/int/None/tensor/etc in tensor[...], got {0!s}.'
class helion.exc.HostTensorDirectUsage(*args, **kwargs)[source]
Parameters:

Raised when host tensors are used directly in device code without proper indexing.

message = "Direct use of host tensor '{0}' in op '{1}' not allowed inside the `hl.tile` or `hl.grid` loop. First load it using {0}[...] or hl.load({0}, ...)."

Assignment and Variable Errors

class helion.exc.RequiresTensorInAssignment(*args, **kwargs)[source]
Parameters:

Raised when non-tensor appears on RHS of assignment.

message = 'Expected tensor in right-hand side of assignment, got {0!s}.'
class helion.exc.NonTensorSubscriptAssign(*args, **kwargs)[source]
Parameters:

Raised for invalid types in subscript assignment.

message = 'Expected tensor in subscript assignment, got {0!s} and {1!s}.'
class helion.exc.AssignmentMultipleTargets(*args, **kwargs)[source]
Parameters:

Raised for multiple assignment targets (a=b=1) on device.

message = 'Assignment with multiple targets (a=b=1) is not allowed inside the `hl.tile` or `hl.grid` loop.'
class helion.exc.InvalidAssignment(*args, **kwargs)[source]
Parameters:

Raised for invalid assignment targets on device.

message = 'Assignment target must be Name or Subscript inside the `hl.tile` or `hl.grid` loop.'
class helion.exc.FailedToUnpackTupleAssign(*args, **kwargs)[source]
Parameters:

Raised when tuple unpacking fails in assignment.

message = 'Failed to unpack values in tuple assignment. Expected a sequence of size {0}, got type: {1!s}.'
class helion.exc.ShapeMismatch(*args, **kwargs)[source]
Parameters:

Raised for shape incompatibility between tensors.

message = 'Shape mismatch between {0!s} and {1!s}.'
class helion.exc.UndefinedVariable(*args, **kwargs)[source]
Parameters:

Raised when referencing undefined variables.

message = '{} is not defined.'
class helion.exc.CannotModifyHostVariableOnDevice(*args, **kwargs)[source]
Parameters:

Raised when modifying host variables inside device loops without subscript assignment.

message = "Cannot modify host variable '{0}' inside `hl.tile` or `hl.grid` loop without subscript assignment. Use '{0}[tile] = ...' instead."
class helion.exc.CannotReadDeviceVariableOnHost(*args, **kwargs)[source]
Parameters:

Raised when attempting to read variables defined inside device loops from host context.

message = "Cannot read variable '{0}' defined inside `hl.tile` or `hl.grid` loop from host code."
class helion.exc.DeviceTensorSubscriptAssignmentNotAllowed(*args, **kwargs)[source]
Parameters:

Raised when attempting to assign to subscript of device tensor.

message = "Cannot assign to subscript of device tensor '{0}'."

Type and Inference Errors

class helion.exc.UnsupportedPythonType(*args, **kwargs)[source]
Parameters:

Raised for Python types not supported in kernels.

message = '{0} is not supported in Helion kernels'
class helion.exc.TypeInferenceError(*args, **kwargs)[source]
Parameters:

Raised for type inference failures.

message = '{0}'
class helion.exc.CantCombineTypesInControlFlow(*args, **kwargs)[source]
Parameters:

Raised for type conflicts in control flow.

message = 'Cannot combine types for {0!r} in control flow: {1} and {2}'
class helion.exc.TracedArgNotSupported(*args, **kwargs)[source]
Parameters:

Raised for unsupported argument types in traced functions.

message = '{!s} is not supported as an arg to traced functions.'

Configuration Errors

class helion.exc.InvalidConfig(*args, **kwargs)[source]
Parameters:

Raised for invalid kernel configurations.

message = '{}'
class helion.exc.NotEnoughConfigs(*args, **kwargs)[source]
Parameters:

Raised when insufficient configs provided for FiniteSearch.

message = 'FiniteSearch requires at least two configs, but got {0}.'
class helion.exc.ShapeSpecializingCall(*args, **kwargs)[source]
Parameters:

Raised for calls requiring shape specialization.

message = 'Call would force shape specialization, try `hl.specialize(x)` or `hl.constexpr`.'
class helion.exc.ShapeSpecializingAllocation(*args, **kwargs)[source]
Parameters:

Raised for allocations requiring specialization.

message = 'Using a tensor size in a device allocation requires specialization. Use `hl.specialize` or `hl.constexpr` to specialize the size.'
class helion.exc.SpecializeOnDevice(*args, **kwargs)[source]
Parameters:

Raised when hl.specialize() is called in device loop.

message = 'hl.specialize() must be called outside the `hl.tile` or `hl.grid` loop.'
class helion.exc.SpecializeArgType(*args, **kwargs)[source]
Parameters:

Raised for invalid arguments to hl.specialize().

message = 'hl.specialize() must be called on a size from an input tensor, got: {}'
class helion.exc.ConfigSpecFragmentWithSymInt(*args, **kwargs)[source]
Parameters:

Raised for ConfigSpecFragment with SymInt.

message = 'ConfigSpecFragment with SymInt arg is not supported. hl.constexpr or hl.specialize may be used to specialize the SymInt value.'
class helion.exc.AutotuningDisallowedInEnvironment(*args, **kwargs)[source]
Parameters:

Raised when autotuning is disabled in environment and no config is provided.

message = 'Autotuning is disabled {0}, please provide a config to @helion.kernel via the config= argument.'

Tunable Parameter Errors

class helion.exc.RegisterTunableArgTypes(*args, **kwargs)[source]
Parameters:

Raised for invalid argument types to hl.register_tunable().

message = 'Expected string literal and ConfigSpecFragment literal, got {0} and {1}.'
class helion.exc.TunableTypeNotSupported(*args, **kwargs)[source]
Parameters:

Raised for unsupported tunable parameter types.

message = 'hl.register_tunable() only supports integer, float, and boolean types, got {0!s}.'
class helion.exc.TunableNameConflict(*args, **kwargs)[source]
Parameters:

Raised for duplicate tunable parameter names.

message = 'Tunable parameter with name {0!s} already exists. Please use a different name.'

Language and Syntax Errors

class helion.exc.NamingConflict(*args, **kwargs)[source]
Parameters:

Raised when reserved variable names are used.

message = 'The variable name {} is reserved in Helion and cannot be used.'
class helion.exc.ClosuresNotSupported(*args, **kwargs)[source]
Parameters:

Raised when closures are found in kernels.

message = 'A closure ({0!r}) was found in the kernel. Closures are not supported.'
class helion.exc.ClosureMutation(*args, **kwargs)[source]
Parameters:

Raised for closure variable mutation.

message = 'Closure mutation (of {0}) is not allowed in a function arg.'
class helion.exc.GlobalMutation(*args, **kwargs)[source]
Parameters:

Raised for global variable mutation.

message = 'Global mutation (of {0}) is not allowed in a function arg.'
class helion.exc.StatementNotSupported(*args, **kwargs)[source]
Parameters:

Raised for unsupported statement types.

message = 'The statement {} is not supported.'
class helion.exc.StarredArgsNotSupportedOnDevice(*args, **kwargs)[source]
Parameters:

Raised for */** args in device loops.

message = '*/** args are not supported inside the `hl.tile` or `hl.grid` loop.'
class helion.exc.DecoratorAfterHelionKernelDecorator(*args, **kwargs)[source]
Parameters:

Raised when decorators appear after @helion.kernel.

message = 'Decorators after helion kernel decorator are not allowed.'

Grid and Execution Errors

class helion.exc.NoTensorArgs(*args, **kwargs)[source]
Parameters:

Raised for kernels with no tensor arguments.

message = 'Kernel took no tensor args, unclear what device to use.'

Compilation and Runtime Errors

class helion.exc.ErrorCompilingKernel(*args, **kwargs)[source]
Parameters:

Raised for compilation failures with error/warning counts.

message = '{0} errors and {1} warnings occurred (see above)'
class helion.exc.TritonError(*args, **kwargs)[source]
Parameters:

Raised for errors in generated Triton programs.

message = 'Error running generated Triton program:\n{1}\n{0}'
class helion.exc.InductorLoweringError(*args, **kwargs)[source]
Parameters:

Raised for Inductor lowering failures.

message = '{}'
class helion.exc.TorchOpTracingError(e)[source]
Parameters:

e (Exception)

Raised for Torch operation tracing errors.

class helion.exc.InternalError(e)[source]
Parameters:

e (Exception)

Raised for internal compiler errors.

Warning Classes

Warnings can be suppressed by including them in the ignore_warnings setting:

class helion.exc.TensorOperationInWrapper(*args, **kwargs)[source]
Parameters:

Warns when tensor operations occur outside hl.tile/hl.grid loops.

message = 'A tensor operation outside of the `hl.tile` or `hl.grid` loop will not be fused in the generated kernel.\nUse @helion.kernel(ignore_warnings=[helion.exc.TensorOperationInWrapper]) to suppress this warning.\nIf this is not a tensor operation, please report this as a bug.'
class helion.exc.TensorOperationsInHostCall(*args, **kwargs)[source]
Parameters:

Specific variant for tensor operations in host calls.

message = 'A tensor operation outside of the `hl.tile` or `hl.grid` loop will not be fused in the generated kernel: {}'
class helion.exc.WrongDevice(*args, **kwargs)[source]
Parameters:

Warns when operations return tensors on wrong device.

message = 'Operation {0} returned a tensor on {1} device, but the kernel is on {2} device.'

Warning Suppression

# Suppress specific warnings
@helion.kernel(ignore_warnings=[helion.exc.TensorOperationInWrapper])
def my_kernel(x):
    y = x * 2  # This operation won't be fused, but warning suppressed
    for i in hl.grid(x.size(0)):
        pass


# Suppress ALL warnings by using BaseWarning
@helion.kernel(ignore_warnings=[helion.exc.BaseWarning])
def quiet_kernel(x):
    # This kernel will suppress all Helion warnings
    pass

See Also

  • Settings - Configuring debug output and warning suppression

  • Kernel - Kernel execution and error contexts