fix(ci): cudaRoundMode typing failure in FP8 test

[kaeun97]

Attempt to fix this issue.

cuda-bindings 13.2.0 changed cudaRoundMode from a standard Python IntEnum to a FastEnumMetaclass type. Numba's type inference cannot resolve FastEnumMetaclass types, causing three FP8 tests to fail (ref).

This PR replaces adds local IntEnum in cuda_fp8.py, matching the pattern already used for saturation_t and `fp8_interpretation_t.

If this works, let file a issue on the Numbast side (if needed).

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[kaeun97]

Thanks in advance - unable to reproduce the error locally so would have to use CI to test this. @gmarkall

[leofang]

@gmarkall @isVoid what does it take for numba-cuda to understand cudaRoundMode from cuda-bindings as a legal enum, without numba-cuda having to repeat the definitions?

On the cuda-bindings side, we switched to a custom fast enum from the builtin IntEnum to reduce Python overhead (cc @mdboom), and we can evaluate if a patch on cuda-bindings makes better sense, or if we could just register the type in numba-cuda (or both).

[leofang]

if it only needs a simple patch to cuda-bindings (and @mdboom deems it's acceptable), it'd be better to fix it on the cuda-bindings side and publish 13.2.1/12.9.7, instead of breaking numba-cuda.

cc @rparolin for vis

[gmarkall]

We'd need to add typing to Numba for the FastEnumMetaclass type for Numba to recognise it.

[leofang]

Would it help if we add __int__ to the fast enum class?

[gmarkall]

I don't think Numba-CUDA will see that. Apart from a small set of cases:

numba-cuda/numba_cuda/numba/cuda/typing/typeof.py

Lines 50 to 97 in 893dde7

	@singledispatch
	def typeof_impl(val, c):
	"""
	Generic typeof() implementation.
	"""
	tp = getattr(val, "_numba_type_", None)
	if tp is not None:
	return tp
	# Check for dlpack objects
	dlpack = getattr(val, "__dlpack__", None)
	if dlpack is not None:
	tp = _typeof_dlpack(dlpack, c)
	if tp is not None:
	return tp
	# Check for __cuda_array_interface__ objects (third-party device arrays)
	# Numba's own DeviceNDArray is handled above via _numba_type_.
	cai = getattr(val, "__cuda_array_interface__", None)
	if cai is not None:
	tp = _typeof_cuda_array_interface(cai, c)
	if tp is not None:
	return tp
	tp = _typeof_buffer(val, c)
	if tp is not None:
	return tp
	# cffi is handled here as it does not expose a public base class
	# for exported functions or CompiledFFI instances.
	from numba.cuda.typing import cffi_utils
	if cffi_utils.SUPPORTED:
	if cffi_utils.is_cffi_func(val):
	return cffi_utils.make_function_type(val)
	if cffi_utils.is_ffi_instance(val):
	return types.ffi
	if HAS_NUMBA:
	# Fallback to Numba's typeof_impl for third-party registrations
	from numba.core.typing.typeof import typeof_impl as core_typeof_impl
	tp = core_typeof_impl(val, c)
	if tp is not None:
	return tp
	return None

Numba-CUDA looks at the Python type to map to the Numba type.

Maybe we could / should add recognition of __int__ as well, but it would need a Numba-CUDA change in addition to adding __int__ to the fast enum class.

[leofang]

Ah, it's done by the same function I just touched recently... Got it. So this is a corner case where Python types do matter and duck-typing does not work.

I would suggest that we merge this PR as a workaround to unblock the CI, and discuss a long-term fix. (My 2c is if the surface area grows to other enums we should register the fast enum type, but one-off patches like this are not bad.) WDYT?

[isVoid]

Sorry I got around to this a bit late. I saw this error yesterday in my Numbast CI but didn't trace it to the bottom. The cuda_fp8.py file was auto generated by Numbast. If we decided on a fast enum type for cuda types we should probably inform Numbast to generate corresponding logics.

[leofang]

/ok to test cfc3787