Repository navigation
[BUG] CI is blocked after cuda-bindings 13.2.0 #833
Description
It seems like CI is blocked after the new release.
The most recent PR that ran CI successfully was this. However, CI fails in this PR and the difference seems to be cuda-bindings version (the failed tests use 13.2.0).
Here is the test failure:
=================================== FAILURES ===================================
_________ FP8Storage_CVT_Intrinsics_Tests.test_cvt_bfloat16raw_to_e8m0 _________
[gw3] linux -- Python 3.12.13 /opt/hostedtoolcache/Python/3.12.13/x64/bin/python
self = <numba.cuda.tests.cudapy.test_fp8_bindings.FP8Storage_CVT_Intrinsics_Tests testMethod=test_cvt_bfloat16raw_to_e8m0>
def test_cvt_bfloat16raw_to_e8m0(self):
@cuda.jit
def kernel(result, exponent):
# NOTE: We currently don't have a good way to construct/modify a
# bfloat16_raw value directly in kernels, so we skip SATFINITE vs
# NOSAT behavior testing for bf16raw here.
raw = cvt_e8m0_to_bf16raw(exponent[0])
result[0] = cvt_bfloat16raw_to_e8m0(
raw, saturation_t.NOSAT, cudaRoundMode.cudaRoundZero
)
result[1] = cvt_bfloat16raw_to_e8m0(
raw, saturation_t.NOSAT, cudaRoundMode.cudaRoundPosInf
)
exponent = np.array([126], dtype=np.uint8) # 0.5
result = np.zeros(2, dtype=np.uint8)
> kernel[1, 1](result, exponent)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py:557:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:719: in __call__
return self.dispatcher.call(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1642: in call
kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1650: in _compile_for_args
return self.compile(tuple(argtypes))
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1908: in compile
kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:155: in __init__
cres = compile_cuda(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/compiler.py:789: in compile_cuda
cres = compile_extra(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/compiler.py:599: in compile_extra
return pipeline.compile_extra(func)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:146: in compile_extra
return self._compile_bytecode()
^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:214: in _compile_bytecode
return self._compile_core()
^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:193: in _compile_core
raise e
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:182: in _compile_core
pm.run(self.state)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:398: in run
raise patched_exception
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:385: in run
self._runPass(idx, pass_inst, state)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:337: in _runPass
mutated |= check(pss.run_pass, internal_state)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:291: in check
mangled = func(compiler_state)
^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typed_passes.py:141: in run_pass
typemap, return_type, calltypes, errs = type_inference_stage(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typed_passes.py:118: in type_inference_stage
infer.build_constraint()
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1108: in build_constraint
self.constrain_statement(inst)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1501: in constrain_statement
self.typeof_assign(inst)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1597: in typeof_assign
self.typeof_global(inst, inst.target, value)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1699: in typeof_global
typ = self.resolve_value_type(inst, gvar.value)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
self = <numba.cuda.core.typeinfer.TypeInferer object at 0x76f4b3965820>
inst = <numba.cuda.core.ir.Assign object at 0x76f4b39572f0>
val = <enum 'cudaRoundMode'>
def resolve_value_type(self, inst, val):
"""
Resolve the type of a simple Python value, such as can be
represented by literals.
"""
try:
return self.context.resolve_value_type(val)
except ValueError as e:
msg = str(e)
> raise TypingError(msg, loc=inst.loc)
E numba.cuda.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
E Untyped global name 'cudaRoundMode': Cannot determine Numba type of <class 'cuda.bindings._internal._fast_enum.FastEnumMetaclass'>
E
E File "../../../../opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py", line 549:
E def kernel(result, exponent):
E <source elided>
E result[0] = cvt_bfloat16raw_to_e8m0(
E raw, saturation_t.NOSAT, cudaRoundMode.cudaRoundZero
E ^
E
E During: Pass nopython_type_inference
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1620: TypingError
___________ FP8Storage_CVT_Intrinsics_Tests.test_cvt_double_to_e8m0 ____________
[gw3] linux -- Python 3.12.13 /opt/hostedtoolcache/Python/3.12.13/x64/bin/python
self = <numba.cuda.tests.cudapy.test_fp8_bindings.FP8Storage_CVT_Intrinsics_Tests testMethod=test_cvt_double_to_e8m0>
def test_cvt_double_to_e8m0(self):
@cuda.jit
def kernel(result, x):
result[0] = cvt_double_to_e8m0(
x[0], saturation_t.NOSAT, cudaRoundMode.cudaRoundPosInf
)
result[1] = cvt_double_to_e8m0(
x[0], saturation_t.SATFINITE, cudaRoundMode.cudaRoundPosInf
)
x_over = np.nextafter(np.float64(2.0**127), np.float64(np.inf))
x = np.array([x_over], dtype=np.float64)
result = np.zeros(2, dtype=np.uint8)
> kernel[1, 1](result, x)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py:534:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:719: in __call__
return self.dispatcher.call(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1642: in call
kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1650: in _compile_for_args
return self.compile(tuple(argtypes))
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1908: in compile
kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:155: in __init__
cres = compile_cuda(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/compiler.py:789: in compile_cuda
cres = compile_extra(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/compiler.py:599: in compile_extra
return pipeline.compile_extra(func)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:146: in compile_extra
return self._compile_bytecode()
^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:214: in _compile_bytecode
return self._compile_core()
^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:193: in _compile_core
raise e
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:182: in _compile_core
pm.run(self.state)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:398: in run
raise patched_exception
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:385: in run
self._runPass(idx, pass_inst, state)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:337: in _runPass
mutated |= check(pss.run_pass, internal_state)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:291: in check
mangled = func(compiler_state)
^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typed_passes.py:141: in run_pass
typemap, return_type, calltypes, errs = type_inference_stage(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typed_passes.py:118: in type_inference_stage
infer.build_constraint()
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1108: in build_constraint
self.constrain_statement(inst)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1501: in constrain_statement
self.typeof_assign(inst)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1597: in typeof_assign
self.typeof_global(inst, inst.target, value)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1699: in typeof_global
typ = self.resolve_value_type(inst, gvar.value)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
self = <numba.cuda.core.typeinfer.TypeInferer object at 0x76f4b5d68cb0>
inst = <numba.cuda.core.ir.Assign object at 0x76f4b5d31940>
val = <enum 'cudaRoundMode'>
def resolve_value_type(self, inst, val):
"""
Resolve the type of a simple Python value, such as can be
represented by literals.
"""
try:
return self.context.resolve_value_type(val)
except ValueError as e:
msg = str(e)
> raise TypingError(msg, loc=inst.loc)
E numba.cuda.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
E Untyped global name 'cudaRoundMode': Cannot determine Numba type of <class 'cuda.bindings._internal._fast_enum.FastEnumMetaclass'>
E
E File "../../../../opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py", line 525:
E def kernel(result, x):
E <source elided>
E result[0] = cvt_double_to_e8m0(
E x[0], saturation_t.NOSAT, cudaRoundMode.cudaRoundPosInf
E ^
E
E During: Pass nopython_type_inference
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1620: TypingError
____________ FP8Storage_CVT_Intrinsics_Tests.test_cvt_float_to_e8m0 ____________
[gw3] linux -- Python 3.12.13 /opt/hostedtoolcache/Python/3.12.13/x64/bin/python
self = <numba.cuda.tests.cudapy.test_fp8_bindings.FP8Storage_CVT_Intrinsics_Tests testMethod=test_cvt_float_to_e8m0>
def test_cvt_float_to_e8m0(self):
@cuda.jit
def kernel(result, x):
# Use a value slightly larger than 2^127 and round up so the
# rounded scale would overflow. Then SATFINITE clips to 2^127 while
# NOSAT produces NaN (0xFF).
result[0] = cvt_float_to_e8m0(
x[0], saturation_t.NOSAT, cudaRoundMode.cudaRoundPosInf
)
result[1] = cvt_float_to_e8m0(
x[0], saturation_t.SATFINITE, cudaRoundMode.cudaRoundPosInf
)
x_over = np.nextafter(np.float32(2.0**127), np.float32(np.inf))
x = np.array([x_over], dtype=np.float32)
result = np.zeros(2, dtype=np.uint8)
> kernel[1, 1](result, x)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py:514:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:719: in __call__
return self.dispatcher.call(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1642: in call
kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1650: in _compile_for_args
return self.compile(tuple(argtypes))
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:1908: in compile
kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/dispatcher.py:155: in __init__
cres = compile_cuda(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/compiler.py:789: in compile_cuda
cres = compile_extra(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/compiler.py:599: in compile_extra
return pipeline.compile_extra(func)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:146: in compile_extra
return self._compile_bytecode()
^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:214: in _compile_bytecode
return self._compile_core()
^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:193: in _compile_core
raise e
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler.py:182: in _compile_core
pm.run(self.state)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:398: in run
raise patched_exception
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:385: in run
self._runPass(idx, pass_inst, state)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_lock.py:74: in _acquire_compile_lock
return func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:337: in _runPass
mutated |= check(pss.run_pass, internal_state)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/compiler_machinery.py:291: in check
mangled = func(compiler_state)
^^^^^^^^^^^^^^^^^^^^
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typed_passes.py:141: in run_pass
typemap, return_type, calltypes, errs = type_inference_stage(
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typed_passes.py:118: in type_inference_stage
infer.build_constraint()
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1108: in build_constraint
self.constrain_statement(inst)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1501: in constrain_statement
self.typeof_assign(inst)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1597: in typeof_assign
self.typeof_global(inst, inst.target, value)
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1699: in typeof_global
typ = self.resolve_value_type(inst, gvar.value)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
self = <numba.cuda.core.typeinfer.TypeInferer object at 0x76f4bc7f76b0>
inst = <numba.cuda.core.ir.Assign object at 0x76f4bc7f65a0>
val = <enum 'cudaRoundMode'>
def resolve_value_type(self, inst, val):
"""
Resolve the type of a simple Python value, such as can be
represented by literals.
"""
try:
return self.context.resolve_value_type(val)
except ValueError as e:
msg = str(e)
> raise TypingError(msg, loc=inst.loc)
E numba.cuda.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
E Untyped global name 'cudaRoundMode': Cannot determine Numba type of <class 'cuda.bindings._internal._fast_enum.FastEnumMetaclass'>
E
E File "../../../../opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py", line 505:
E def kernel(result, x):
E <source elided>
E result[0] = cvt_float_to_e8m0(
E x[0], saturation_t.NOSAT, cudaRoundMode.cudaRoundPosInf
E ^
E
E During: Pass nopython_type_inference
/opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/core/typeinfer.py:1620: TypingError
=========================== short test summary info ============================
FAILED cudapy/test_fp8_bindings.py::FP8Storage_CVT_Intrinsics_Tests::test_cvt_bfloat16raw_to_e8m0 - numba.cuda.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
Untyped global name 'cudaRoundMode': Cannot determine Numba type of <class 'cuda.bindings._internal._fast_enum.FastEnumMetaclass'>
File "../../../../opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py", line 549:
def kernel(result, exponent):
<source elided>
result[0] = cvt_bfloat16raw_to_e8m0(
raw, saturation_t.NOSAT, cudaRoundMode.cudaRoundZero
^
During: Pass nopython_type_inference
FAILED cudapy/test_fp8_bindings.py::FP8Storage_CVT_Intrinsics_Tests::test_cvt_double_to_e8m0 - numba.cuda.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
Untyped global name 'cudaRoundMode': Cannot determine Numba type of <class 'cuda.bindings._internal._fast_enum.FastEnumMetaclass'>
File "../../../../opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py", line 525:
def kernel(result, x):
<source elided>
result[0] = cvt_double_to_e8m0(
x[0], saturation_t.NOSAT, cudaRoundMode.cudaRoundPosInf
^
During: Pass nopython_type_inference
FAILED cudapy/test_fp8_bindings.py::FP8Storage_CVT_Intrinsics_Tests::test_cvt_float_to_e8m0 - numba.cuda.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
Untyped global name 'cudaRoundMode': Cannot determine Numba type of <class 'cuda.bindings._internal._fast_enum.FastEnumMetaclass'>
File "../../../../opt/hostedtoolcache/Python/3.12.13/x64/lib/python3.12/site-packages/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py", line 505:
def kernel(result, x):
<source elided>
result[0] = cvt_float_to_e8m0(
x[0], saturation_t.NOSAT, cudaRoundMode.cudaRoundPosInf
^
During: Pass nopython_type_inference
====== 3 failed, 1970 passed, 44 skipped, 13 xfailed in 75.34s (0:01:15) =======
Reacted by Leo Fang
Activity
fix(ci): cudaRoundMode typing failure in FP8 test (#834)
gmarkall commentedon Mar 12, 2026
@kaeun97 Many thanks for the PR that unblocks CI! Would you be interested in making a PR that adds support for CUDA Python's
FastIntEnumMetaclassto Numba-CUDA? (no worries if not, just let me know 🙂 )kaeun97 commentedon Mar 12, 2026
Of course, thank you @gmarkall ! Will update you with a PR :)
mdboom commentedon Mar 13, 2026
@kauen97: I'm close to having a workaround in cuda_bindings for this that won't require any changes to numba-cuda or numba. Will post soon.
kaeun97 commentedon Mar 13, 2026
Awesome @mdboom , I was planning to submit an issue there, but glad that you are already working on it. Thank you!
mdboom commentedon Mar 13, 2026
Actually, I was unable to find a way to support this externally. However, I'd like to propose an alternative solution in #837 which won't require any enums shared between numba-core and cuda-bindings to be kept in sync.
rparolin commentedon Mar 23, 2026
Fixed by #837.