diff --git a/configs/cuda_fp8.yml b/configs/cuda_fp8.yml new file mode 100644 index 000000000..c18464c32 --- /dev/null +++ b/configs/cuda_fp8.yml @@ -0,0 +1,38 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause +Name: Numba FP8 +Version: 0.0.1 +GPU Arch: + - sm_90 +Entry Point: ./numba_cuda/numba/cuda/include/13/cuda_fp8.h +Clang Include Paths: + - ./numba_cuda/numba/cuda/include/13 +File List: + - ./numba_cuda/numba/cuda/include/13/cuda_fp8.h + - ./numba_cuda/numba/cuda/include/13/cuda_fp8.hpp +Exclude: {} +Types: + __nv_fp8_e5m2: Number + __nv_fp8_e4m3: Number + __nv_fp8_e8m0: Number + __nv_fp8x2_e5m2: Type + __nv_fp8x2_e4m3: Type + __nv_fp8x2_e8m0: Type +Data Models: + __nv_fp8_e5m2: PrimitiveModel + __nv_fp8_e4m3: PrimitiveModel + __nv_fp8_e8m0: PrimitiveModel + __nv_fp8x2_e5m2: StructModel + __nv_fp8x2_e4m3: StructModel + __nv_fp8x2_e8m0: StructModel +API Prefix Removal: + Function: + - __nv_ + Struct: + - __nv_ + Enum: + - __nv_ + - __NV_ +Skip Prefix: "__internal" +Shim Include Override: "\"cuda_fp8.h\"" +Use Separate Registry: True diff --git a/numba_cuda/numba/cuda/_internal/cuda_fp8.py b/numba_cuda/numba/cuda/_internal/cuda_fp8.py new file mode 100644 index 000000000..2f255fe32 --- /dev/null +++ b/numba_cuda/numba/cuda/_internal/cuda_fp8.py @@ -0,0 +1,3936 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# Automatically generated by Numbast Static Binding Generator +# Generator Information: +# Ast_canopy version: 0.6.1.dev18+gc41f52f6f.d20251223 +# Numbast version: 0.6.1.dev0+ga900697e9.d20251217 +# Generation command: static_binding_generator --cfg-path configs/cuda_fp8.yml --output-dir ./numba_cuda/numba/cuda/_internal +# Static binding generator parameters: {'cfg_path': 'configs/cuda_fp8.yml', 'output_dir': './numba_cuda/numba/cuda/_internal', 'run_ruff_format': True, 'bypass_parse_error': False} +# Config file path (relative to the path of the generated binding): ../../../../../configs/cuda_fp8.yml +# Cudatoolkit version: (13, 0) + + +# Imports: +from numba.cuda.typing.templates import ConcreteTemplate +from numba.cuda.types import float16 +from numba.cuda.types import IntEnumMember +from numba import types +from numba.cuda.extending import make_attribute_wrapper +from numba.cuda.types import bool_ +from numba.cuda.types import uint64 +from llvmlite import ir +from cuda.bindings.runtime import cudaRoundMode +from numba.cuda.typing.templates import Registry as TypingRegistry +from numba.cuda.types import uint32 +from numba.cuda.types import float32 +from enum import IntEnum +from numba.cuda.datamodel import StructModel +from numba.cuda.core.imputils import Registry as TargetRegistry +from numba.cuda.extending import as_numba_type +from numba.cuda.types import Function +from numba.cuda.types import Type +from numba.cuda.types import int32 +from numba.types import BoundFunction +from numba.cuda.types import uint8 +from numba.cuda.typing.templates import AttributeTemplate +from numba.cuda.vector_types import vector_types +from numba.cuda.types import int8 +from numba.core.typing import signature +from numba.cuda.types import bfloat16 +from numba.cuda.typing import signature +from numba.cuda.extending import register_model +from numba.cuda.types import float64 +from numba.cuda.datamodel import PrimitiveModel +from numba.cuda._internal.cuda_bf16 import ( + _type_unnamed1405307 as bfloat16_raw_type, +) +import numba +from numba.cuda.types import int16 +from numba.cuda.types import int64 +import io +from numba.cuda.types import Number +from numba.cuda.types import uint16 +from numba.cuda import CUSource + + +float32x2 = vector_types["float32x2"] +float64x2 = vector_types["float64x2"] +float32x4 = vector_types["float32x4"] +float64x4 = vector_types["float64x4"] + + +typing_registry = TypingRegistry() +register = typing_registry.register +register_attr = typing_registry.register_attr +register_global = typing_registry.register_global +target_registry = TargetRegistry() +lower = target_registry.lower +lower_attr = target_registry.lower_getattr +lower_constant = target_registry.lower_constant +lower_cast = target_registry.lower_cast + +# Shim Stream: + + +class _KeyedStringIO(io.StringIO): + def __init__(self, *arg, **kwarg): + super().__init__(*arg, *kwarg) + self._keys = set() + + def write_with_key(self, key: str, value: str): + if key in self._keys: + return + self._keys.add(key) + self.write(value) + + def reset(self): + self._keys.clear() + self.seek(0) + + +shim_defines = "" +shim_include = "#include <" + "cuda_fp8.h" + ">" +shim_prefix = shim_defines + "\n" + shim_include +shim_stream = _KeyedStringIO() +shim_stream.write(shim_prefix) +shim_obj = CUSource(shim_stream) + + +from numba.cuda.target import CUDATargetContext + +from llvmlite import ir + + +def prepare_ir_types( + context: CUDATargetContext, argtys: list[ir.Type] +) -> list[ir.Type]: + """ + Prepare IR types for passing arguments via pointers in function calls. + + This utility wraps each argument type in a PointerType to enable + the call convention used by FunctionCallConv, where arguments are + passed by reference. + + Parameters + ---------- + context : context object + The compilation context providing the get_value_type method. + argtys : list[ir.Type] + List of LLVM IR types representing function arguments. + + Returns + ------- + list[ir.Type] + List of pointer types wrapping the value types of each argument. + """ + return [ir.PointerType(context.get_value_type(argty)) for argty in argtys] + + +from numba.cuda import types, cgutils + +from llvmlite import ir + + +class BaseCallConv: + shim_function_template = "{mangled_name}_nbst" + + def __init__( + self, + itanium_mangled_name: str, + shim_writer: object, + shim_code: str, + ): + self.shim_writer = shim_writer + self.itanium_mangled_name = itanium_mangled_name + self.shim_code = shim_code + + self.shim_function_name = self.shim_function_template.format( + mangled_name=self.itanium_mangled_name + ) + + def _lazy_write_shim(self, shim_code: str): + self.shim_writer.write_to_shim(shim_code, self.shim_function_name) + + def _lower(self, builder, context, sig, args): + self._lazy_write_shim(self.shim_code) + return self._lower_impl(builder, context, sig, args) + + def _lower_impl(self, builder, context, sig, args): + raise NotImplementedError + + def __call__(self, builder, context, sig, args): + return self._lower(builder, context, sig, args) + + +class FunctionCallConv(BaseCallConv): + def _lower_impl(self, builder, context, sig, args): + return_type = sig.return_type + # 1. Prepare return value pointer + if return_type == types.void: + # Void return type in C++ is shimmed as int& ignored + retval_ty = ir.IntType(32) + retval_ptr = builder.alloca(retval_ty, name="ignored") + else: + retval_ty = context.get_value_type(return_type) + retval_ptr = builder.alloca(retval_ty, name="retval") + + # 2. Prepare arguments + arg_pointer_types = prepare_ir_types(context, sig.args) + + # All arguments are passed by pointer + ptrs = [ + cgutils.alloca_once(builder, context.get_value_type(argty)) + for argty in sig.args + ] + for ptr, argty, arg in zip(ptrs, sig.args, args): + builder.store(arg, ptr, align=getattr(argty, "alignof_", None)) + + # 3. Declare shim + # Shim signature: int (retval_type*, arg0_type*, ...) + fnty = ir.FunctionType( + ir.IntType(32), [ir.PointerType(retval_ty)] + arg_pointer_types + ) + fn = cgutils.get_or_insert_function( + builder.module, fnty, self.shim_function_name + ) + + # 4. Call shim + builder.call(fn, (retval_ptr, *ptrs)) + + # 5. Return + if return_type == types.void: + return None + else: + return builder.load( + retval_ptr, align=getattr(return_type, "alignof_", None) + ) + + +class ShimWriterAdapter: + def __init__(self, stream): + self.stream = stream + + def write_to_shim(self, content, id): + self.stream.write_with_key(id, content) + + +shim_writer = ShimWriterAdapter(shim_stream) + +# Enums: + + +class saturation_t(IntEnum): + NOSAT = 0 + SATFINITE = 1 + + +class fp8_interpretation_t(IntEnum): + E4M3 = 0 + E5M2 = 1 + + +# Structs: + + +# Typing for fp8_e5m2 +class _type_class_fp8_e5m2(Number): + def __init__(self): + super().__init__(name="fp8_e5m2") + self.alignof_ = 1 + self.bitwidth = 1 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8_e5m2 = _type_class_fp8_e5m2() + + +# Make Python API for struct +fp8_e5m2 = type("fp8_e5m2", (), {"_nbtype": _type_fp8_e5m2}) + +as_numba_type.register(fp8_e5m2, _type_fp8_e5m2) + + +@register_model(_type_class_fp8_e5m2) +class _model_fp8_e5m2(PrimitiveModel): + def __init__(self, dmm, fe_type): + be_type = ir.IntType(fe_type.bitwidth) + super(_model_fp8_e5m2, self).__init__(dmm, fe_type, be_type) + + +def _lower__ZN13__nv_fp8_e5m2C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Ev_nbst(__nv_fp8_e5m2 *self ) { + new (self) __nv_fp8_e5m2(); + return 0; + } + """ + + @lower( + fp8_e5m2, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Ev(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1E6__half(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1E6__half_nbst(__nv_fp8_e5m2 *self , __half* f) { + new (self) __nv_fp8_e5m2(*f); + return 0; + } + """ + + @lower(fp8_e5m2, float16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1E6__half", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1E6__half(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1E13__nv_bfloat16(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1E13__nv_bfloat16_nbst(__nv_fp8_e5m2 *self , __nv_bfloat16* f) { + new (self) __nv_fp8_e5m2(*f); + return 0; + } + """ + + @lower(fp8_e5m2, bfloat16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1E13__nv_bfloat16", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1E13__nv_bfloat16(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Ef(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Ef_nbst(__nv_fp8_e5m2 *self , float* f) { + new (self) __nv_fp8_e5m2(*f); + return 0; + } + """ + + @lower(fp8_e5m2, float32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Ef", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Ef(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Ed(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Ed_nbst(__nv_fp8_e5m2 *self , double* f) { + new (self) __nv_fp8_e5m2(*f); + return 0; + } + """ + + @lower(fp8_e5m2, float64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Ed", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Ed(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Et(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Et_nbst(__nv_fp8_e5m2 *self , unsigned short* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, uint16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Et", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Et(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Ej(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Ej_nbst(__nv_fp8_e5m2 *self , unsigned int* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, uint32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Ej", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Ej(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Em(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Em_nbst(__nv_fp8_e5m2 *self , unsigned long* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, uint64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Em", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Em(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Ey(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Ey_nbst(__nv_fp8_e5m2 *self , unsigned long long* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, uint64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Ey", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Ey(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Es(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Es_nbst(__nv_fp8_e5m2 *self , short* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, int16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Es", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Es(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Ei(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Ei_nbst(__nv_fp8_e5m2 *self , int* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, int32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Ei", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Ei(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1El(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1El_nbst(__nv_fp8_e5m2 *self , long* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, int64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1El", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1El(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e5m2C1Ex(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e5m2C1Ex_nbst(__nv_fp8_e5m2 *self , long long* val) { + new (self) __nv_fp8_e5m2(*val); + return 0; + } + """ + + @lower(fp8_e5m2, int64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e5m2C1Ex", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e5m2C1Ex(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8_e5m2(ConcreteTemplate): + key = globals()["fp8_e5m2"] + cases = [ + signature( + _type_fp8_e5m2, + ), + signature(_type_fp8_e5m2, float16), + signature(_type_fp8_e5m2, bfloat16), + signature(_type_fp8_e5m2, float32), + signature(_type_fp8_e5m2, float64), + signature(_type_fp8_e5m2, uint16), + signature(_type_fp8_e5m2, uint32), + signature(_type_fp8_e5m2, uint64), + signature(_type_fp8_e5m2, uint64), + signature(_type_fp8_e5m2, int16), + signature(_type_fp8_e5m2, int32), + signature(_type_fp8_e5m2, int64), + signature(_type_fp8_e5m2, int64), + ] + + +register_global(fp8_e5m2, Function(_ctor_template___nv_fp8_e5m2)) + + +def _from___nv_fp8_e5m2_to_float16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cv6__halfEv_nbst(__half &retval, __nv_fp8_e5m2 *self) { + retval = self->operator __half(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, float16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cv6__halfEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_float16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_float32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvfEv_nbst(float &retval, __nv_fp8_e5m2 *self) { + retval = self->operator float(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, float32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvfEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_float32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_bfloat16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cv13__nv_bfloat16Ev_nbst(__nv_bfloat16 &retval, __nv_fp8_e5m2 *self) { + retval = self->operator __nv_bfloat16(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, bfloat16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cv13__nv_bfloat16Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_bfloat16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_float64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvdEv_nbst(double &retval, __nv_fp8_e5m2 *self) { + retval = self->operator double(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, float64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvdEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_float64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_uint8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvhEv_nbst(unsigned char &retval, __nv_fp8_e5m2 *self) { + retval = self->operator unsigned char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, uint8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvhEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_uint8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_uint16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvtEv_nbst(unsigned short &retval, __nv_fp8_e5m2 *self) { + retval = self->operator unsigned short(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, uint16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvtEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_uint16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_uint32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvjEv_nbst(unsigned int &retval, __nv_fp8_e5m2 *self) { + retval = self->operator unsigned int(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, uint32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvjEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_uint32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_uint64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvmEv_nbst(unsigned long &retval, __nv_fp8_e5m2 *self) { + retval = self->operator unsigned long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, uint64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvmEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_uint64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_int8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvaEv_nbst(signed char &retval, __nv_fp8_e5m2 *self) { + retval = self->operator signed char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, int8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvaEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_int8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_int16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvsEv_nbst(short &retval, __nv_fp8_e5m2 *self) { + retval = self->operator short(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, int16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvsEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_int16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_int32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cviEv_nbst(int &retval, __nv_fp8_e5m2 *self) { + retval = self->operator int(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, int32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cviEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_int32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_int64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvlEv_nbst(long &retval, __nv_fp8_e5m2 *self) { + retval = self->operator long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, int64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvlEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_int64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_int64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvxEv_nbst(long long &retval, __nv_fp8_e5m2 *self) { + retval = self->operator long long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, int64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvxEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_int64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e5m2_to_bool__lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e5m2cvbEv_nbst(bool &retval, __nv_fp8_e5m2 *self) { + retval = self->operator bool(); + return 0; + } + """ + + @lower_cast(_type_fp8_e5m2, bool_) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e5m2cvbEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e5m2_to_bool__lower(shim_stream, shim_obj) + + +# Typing for fp8x2_e5m2 +class _type_class_fp8x2_e5m2(Type): + def __init__(self): + super().__init__(name="fp8x2_e5m2") + self.alignof_ = 2 + self.bitwidth = 2 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8x2_e5m2 = _type_class_fp8x2_e5m2() + + +# Make Python API for struct +fp8x2_e5m2 = type("fp8x2_e5m2", (), {"_nbtype": _type_fp8x2_e5m2}) + +as_numba_type.register(fp8x2_e5m2, _type_fp8x2_e5m2) + + +@register_model(_type_class_fp8x2_e5m2) +class _model_fp8x2_e5m2(StructModel): + def __init__(self, dmm, fe_type): + members = [("__x", uint16)] + super().__init__(dmm, fe_type, members) + + +@register_attr +class _attr_typing_fp8x2_e5m2(AttributeTemplate): + key = _type_fp8x2_e5m2 + + def resolve___x(self, obj): + return uint16 + + +make_attribute_wrapper(_type_class_fp8x2_e5m2, "__x", "__x") + + +def _lower__ZN15__nv_fp8x2_e5m2C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e5m2C1Ev_nbst(__nv_fp8x2_e5m2 *self ) { + new (self) __nv_fp8x2_e5m2(); + return 0; + } + """ + + @lower( + fp8x2_e5m2, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e5m2C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e5m2C1Ev(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x2_e5m2C1E6float2(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e5m2C1E6float2_nbst(__nv_fp8x2_e5m2 *self , float2* f) { + new (self) __nv_fp8x2_e5m2(*f); + return 0; + } + """ + + @lower(fp8x2_e5m2, float32x2) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e5m2C1E6float2", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e5m2C1E6float2(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x2_e5m2C1E7double2(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e5m2C1E7double2_nbst(__nv_fp8x2_e5m2 *self , double2* f) { + new (self) __nv_fp8x2_e5m2(*f); + return 0; + } + """ + + @lower(fp8x2_e5m2, float64x2) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e5m2C1E7double2", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e5m2C1E7double2(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8x2_e5m2(ConcreteTemplate): + key = globals()["fp8x2_e5m2"] + cases = [ + signature( + _type_fp8x2_e5m2, + ), + signature(_type_fp8x2_e5m2, float32x2), + signature(_type_fp8x2_e5m2, float64x2), + ] + + +register_global(fp8x2_e5m2, Function(_ctor_template___nv_fp8x2_e5m2)) + + +def _from___nv_fp8x2_e5m2_to_float32x2_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK15__nv_fp8x2_e5m2cv6float2Ev_nbst(float2 &retval, __nv_fp8x2_e5m2 *self) { + retval = self->operator float2(); + return 0; + } + """ + + @lower_cast(_type_fp8x2_e5m2, float32x2) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK15__nv_fp8x2_e5m2cv6float2Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8x2_e5m2_to_float32x2_lower(shim_stream, shim_obj) + + +# Typing for fp8x4_e5m2 +class _type_class_fp8x4_e5m2(Type): + def __init__(self): + super().__init__(name="fp8x4_e5m2") + self.alignof_ = 4 + self.bitwidth = 4 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8x4_e5m2 = _type_class_fp8x4_e5m2() + + +# Make Python API for struct +fp8x4_e5m2 = type("fp8x4_e5m2", (), {"_nbtype": _type_fp8x4_e5m2}) + +as_numba_type.register(fp8x4_e5m2, _type_fp8x4_e5m2) + + +@register_model(_type_class_fp8x4_e5m2) +class _model_fp8x4_e5m2(StructModel): + def __init__(self, dmm, fe_type): + members = [("__x", uint32)] + super().__init__(dmm, fe_type, members) + + +@register_attr +class _attr_typing_fp8x4_e5m2(AttributeTemplate): + key = _type_fp8x4_e5m2 + + def resolve___x(self, obj): + return uint32 + + +make_attribute_wrapper(_type_class_fp8x4_e5m2, "__x", "__x") + + +def _lower__ZN15__nv_fp8x4_e5m2C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e5m2C1Ev_nbst(__nv_fp8x4_e5m2 *self ) { + new (self) __nv_fp8x4_e5m2(); + return 0; + } + """ + + @lower( + fp8x4_e5m2, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e5m2C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e5m2C1Ev(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x4_e5m2C1E6float4(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e5m2C1E6float4_nbst(__nv_fp8x4_e5m2 *self , float4* f) { + new (self) __nv_fp8x4_e5m2(*f); + return 0; + } + """ + + @lower(fp8x4_e5m2, float32x4) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e5m2C1E6float4", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e5m2C1E6float4(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x4_e5m2C1E7double4(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e5m2C1E7double4_nbst(__nv_fp8x4_e5m2 *self , double4* f) { + new (self) __nv_fp8x4_e5m2(*f); + return 0; + } + """ + + @lower(fp8x4_e5m2, float64x4) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e5m2C1E7double4", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e5m2C1E7double4(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8x4_e5m2(ConcreteTemplate): + key = globals()["fp8x4_e5m2"] + cases = [ + signature( + _type_fp8x4_e5m2, + ), + signature(_type_fp8x4_e5m2, float32x4), + signature(_type_fp8x4_e5m2, float64x4), + ] + + +register_global(fp8x4_e5m2, Function(_ctor_template___nv_fp8x4_e5m2)) + + +def _from___nv_fp8x4_e5m2_to_float32x4_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK15__nv_fp8x4_e5m2cv6float4Ev_nbst(float4 &retval, __nv_fp8x4_e5m2 *self) { + retval = self->operator float4(); + return 0; + } + """ + + @lower_cast(_type_fp8x4_e5m2, float32x4) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK15__nv_fp8x4_e5m2cv6float4Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8x4_e5m2_to_float32x4_lower(shim_stream, shim_obj) + + +# Typing for fp8_e4m3 +class _type_class_fp8_e4m3(Number): + def __init__(self): + super().__init__(name="fp8_e4m3") + self.alignof_ = 1 + self.bitwidth = 1 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8_e4m3 = _type_class_fp8_e4m3() + + +# Make Python API for struct +fp8_e4m3 = type("fp8_e4m3", (), {"_nbtype": _type_fp8_e4m3}) + +as_numba_type.register(fp8_e4m3, _type_fp8_e4m3) + + +@register_model(_type_class_fp8_e4m3) +class _model_fp8_e4m3(PrimitiveModel): + def __init__(self, dmm, fe_type): + be_type = ir.IntType(fe_type.bitwidth) + super(_model_fp8_e4m3, self).__init__(dmm, fe_type, be_type) + + +def _lower__ZN13__nv_fp8_e4m3C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Ev_nbst(__nv_fp8_e4m3 *self ) { + new (self) __nv_fp8_e4m3(); + return 0; + } + """ + + @lower( + fp8_e4m3, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Ev(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1E6__half(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1E6__half_nbst(__nv_fp8_e4m3 *self , __half* f) { + new (self) __nv_fp8_e4m3(*f); + return 0; + } + """ + + @lower(fp8_e4m3, float16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1E6__half", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1E6__half(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1E13__nv_bfloat16(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1E13__nv_bfloat16_nbst(__nv_fp8_e4m3 *self , __nv_bfloat16* f) { + new (self) __nv_fp8_e4m3(*f); + return 0; + } + """ + + @lower(fp8_e4m3, bfloat16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1E13__nv_bfloat16", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1E13__nv_bfloat16(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Ef(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Ef_nbst(__nv_fp8_e4m3 *self , float* f) { + new (self) __nv_fp8_e4m3(*f); + return 0; + } + """ + + @lower(fp8_e4m3, float32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Ef", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Ef(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Ed(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Ed_nbst(__nv_fp8_e4m3 *self , double* f) { + new (self) __nv_fp8_e4m3(*f); + return 0; + } + """ + + @lower(fp8_e4m3, float64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Ed", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Ed(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Et(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Et_nbst(__nv_fp8_e4m3 *self , unsigned short* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, uint16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Et", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Et(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Ej(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Ej_nbst(__nv_fp8_e4m3 *self , unsigned int* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, uint32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Ej", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Ej(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Em(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Em_nbst(__nv_fp8_e4m3 *self , unsigned long* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, uint64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Em", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Em(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Ey(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Ey_nbst(__nv_fp8_e4m3 *self , unsigned long long* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, uint64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Ey", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Ey(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Es(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Es_nbst(__nv_fp8_e4m3 *self , short* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, int16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Es", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Es(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Ei(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Ei_nbst(__nv_fp8_e4m3 *self , int* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, int32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Ei", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Ei(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1El(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1El_nbst(__nv_fp8_e4m3 *self , long* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, int64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1El", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1El(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e4m3C1Ex(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e4m3C1Ex_nbst(__nv_fp8_e4m3 *self , long long* val) { + new (self) __nv_fp8_e4m3(*val); + return 0; + } + """ + + @lower(fp8_e4m3, int64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e4m3C1Ex", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e4m3C1Ex(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8_e4m3(ConcreteTemplate): + key = globals()["fp8_e4m3"] + cases = [ + signature( + _type_fp8_e4m3, + ), + signature(_type_fp8_e4m3, float16), + signature(_type_fp8_e4m3, bfloat16), + signature(_type_fp8_e4m3, float32), + signature(_type_fp8_e4m3, float64), + signature(_type_fp8_e4m3, uint16), + signature(_type_fp8_e4m3, uint32), + signature(_type_fp8_e4m3, uint64), + signature(_type_fp8_e4m3, uint64), + signature(_type_fp8_e4m3, int16), + signature(_type_fp8_e4m3, int32), + signature(_type_fp8_e4m3, int64), + signature(_type_fp8_e4m3, int64), + ] + + +register_global(fp8_e4m3, Function(_ctor_template___nv_fp8_e4m3)) + + +def _from___nv_fp8_e4m3_to_float16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cv6__halfEv_nbst(__half &retval, __nv_fp8_e4m3 *self) { + retval = self->operator __half(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, float16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cv6__halfEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_float16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_float32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvfEv_nbst(float &retval, __nv_fp8_e4m3 *self) { + retval = self->operator float(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, float32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvfEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_float32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_bfloat16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cv13__nv_bfloat16Ev_nbst(__nv_bfloat16 &retval, __nv_fp8_e4m3 *self) { + retval = self->operator __nv_bfloat16(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, bfloat16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cv13__nv_bfloat16Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_bfloat16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_float64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvdEv_nbst(double &retval, __nv_fp8_e4m3 *self) { + retval = self->operator double(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, float64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvdEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_float64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_uint8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvhEv_nbst(unsigned char &retval, __nv_fp8_e4m3 *self) { + retval = self->operator unsigned char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, uint8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvhEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_uint8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_uint16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvtEv_nbst(unsigned short &retval, __nv_fp8_e4m3 *self) { + retval = self->operator unsigned short(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, uint16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvtEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_uint16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_uint32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvjEv_nbst(unsigned int &retval, __nv_fp8_e4m3 *self) { + retval = self->operator unsigned int(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, uint32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvjEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_uint32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_uint64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvmEv_nbst(unsigned long &retval, __nv_fp8_e4m3 *self) { + retval = self->operator unsigned long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, uint64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvmEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_uint64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_uint64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvyEv_nbst(unsigned long long &retval, __nv_fp8_e4m3 *self) { + retval = self->operator unsigned long long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, uint64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvyEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_uint64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_int8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvaEv_nbst(signed char &retval, __nv_fp8_e4m3 *self) { + retval = self->operator signed char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, int8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvaEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_int8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_int8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvcEv_nbst(char &retval, __nv_fp8_e4m3 *self) { + retval = self->operator char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, int8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvcEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_int8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_int16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvsEv_nbst(short &retval, __nv_fp8_e4m3 *self) { + retval = self->operator short(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, int16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvsEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_int16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_int32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cviEv_nbst(int &retval, __nv_fp8_e4m3 *self) { + retval = self->operator int(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, int32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cviEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_int32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_int64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvlEv_nbst(long &retval, __nv_fp8_e4m3 *self) { + retval = self->operator long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, int64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvlEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_int64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_int64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvxEv_nbst(long long &retval, __nv_fp8_e4m3 *self) { + retval = self->operator long long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, int64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvxEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_int64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e4m3_to_bool__lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e4m3cvbEv_nbst(bool &retval, __nv_fp8_e4m3 *self) { + retval = self->operator bool(); + return 0; + } + """ + + @lower_cast(_type_fp8_e4m3, bool_) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e4m3cvbEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e4m3_to_bool__lower(shim_stream, shim_obj) + + +# Typing for fp8x2_e4m3 +class _type_class_fp8x2_e4m3(Type): + def __init__(self): + super().__init__(name="fp8x2_e4m3") + self.alignof_ = 2 + self.bitwidth = 2 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8x2_e4m3 = _type_class_fp8x2_e4m3() + + +# Make Python API for struct +fp8x2_e4m3 = type("fp8x2_e4m3", (), {"_nbtype": _type_fp8x2_e4m3}) + +as_numba_type.register(fp8x2_e4m3, _type_fp8x2_e4m3) + + +@register_model(_type_class_fp8x2_e4m3) +class _model_fp8x2_e4m3(StructModel): + def __init__(self, dmm, fe_type): + members = [("__x", uint16)] + super().__init__(dmm, fe_type, members) + + +@register_attr +class _attr_typing_fp8x2_e4m3(AttributeTemplate): + key = _type_fp8x2_e4m3 + + def resolve___x(self, obj): + return uint16 + + +make_attribute_wrapper(_type_class_fp8x2_e4m3, "__x", "__x") + + +def _lower__ZN15__nv_fp8x2_e4m3C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e4m3C1Ev_nbst(__nv_fp8x2_e4m3 *self ) { + new (self) __nv_fp8x2_e4m3(); + return 0; + } + """ + + @lower( + fp8x2_e4m3, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e4m3C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e4m3C1Ev(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x2_e4m3C1E6float2(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e4m3C1E6float2_nbst(__nv_fp8x2_e4m3 *self , float2* f) { + new (self) __nv_fp8x2_e4m3(*f); + return 0; + } + """ + + @lower(fp8x2_e4m3, float32x2) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e4m3C1E6float2", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e4m3C1E6float2(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x2_e4m3C1E7double2(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e4m3C1E7double2_nbst(__nv_fp8x2_e4m3 *self , double2* f) { + new (self) __nv_fp8x2_e4m3(*f); + return 0; + } + """ + + @lower(fp8x2_e4m3, float64x2) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e4m3C1E7double2", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e4m3C1E7double2(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8x2_e4m3(ConcreteTemplate): + key = globals()["fp8x2_e4m3"] + cases = [ + signature( + _type_fp8x2_e4m3, + ), + signature(_type_fp8x2_e4m3, float32x2), + signature(_type_fp8x2_e4m3, float64x2), + ] + + +register_global(fp8x2_e4m3, Function(_ctor_template___nv_fp8x2_e4m3)) + + +def _from___nv_fp8x2_e4m3_to_float32x2_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK15__nv_fp8x2_e4m3cv6float2Ev_nbst(float2 &retval, __nv_fp8x2_e4m3 *self) { + retval = self->operator float2(); + return 0; + } + """ + + @lower_cast(_type_fp8x2_e4m3, float32x2) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK15__nv_fp8x2_e4m3cv6float2Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8x2_e4m3_to_float32x2_lower(shim_stream, shim_obj) + + +# Typing for fp8x4_e4m3 +class _type_class_fp8x4_e4m3(Type): + def __init__(self): + super().__init__(name="fp8x4_e4m3") + self.alignof_ = 4 + self.bitwidth = 4 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8x4_e4m3 = _type_class_fp8x4_e4m3() + + +# Make Python API for struct +fp8x4_e4m3 = type("fp8x4_e4m3", (), {"_nbtype": _type_fp8x4_e4m3}) + +as_numba_type.register(fp8x4_e4m3, _type_fp8x4_e4m3) + + +@register_model(_type_class_fp8x4_e4m3) +class _model_fp8x4_e4m3(StructModel): + def __init__(self, dmm, fe_type): + members = [("__x", uint32)] + super().__init__(dmm, fe_type, members) + + +@register_attr +class _attr_typing_fp8x4_e4m3(AttributeTemplate): + key = _type_fp8x4_e4m3 + + def resolve___x(self, obj): + return uint32 + + +make_attribute_wrapper(_type_class_fp8x4_e4m3, "__x", "__x") + + +def _lower__ZN15__nv_fp8x4_e4m3C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e4m3C1Ev_nbst(__nv_fp8x4_e4m3 *self ) { + new (self) __nv_fp8x4_e4m3(); + return 0; + } + """ + + @lower( + fp8x4_e4m3, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e4m3C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e4m3C1Ev(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x4_e4m3C1E6float4(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e4m3C1E6float4_nbst(__nv_fp8x4_e4m3 *self , float4* f) { + new (self) __nv_fp8x4_e4m3(*f); + return 0; + } + """ + + @lower(fp8x4_e4m3, float32x4) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e4m3C1E6float4", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e4m3C1E6float4(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x4_e4m3C1E7double4(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e4m3C1E7double4_nbst(__nv_fp8x4_e4m3 *self , double4* f) { + new (self) __nv_fp8x4_e4m3(*f); + return 0; + } + """ + + @lower(fp8x4_e4m3, float64x4) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e4m3C1E7double4", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e4m3C1E7double4(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8x4_e4m3(ConcreteTemplate): + key = globals()["fp8x4_e4m3"] + cases = [ + signature( + _type_fp8x4_e4m3, + ), + signature(_type_fp8x4_e4m3, float32x4), + signature(_type_fp8x4_e4m3, float64x4), + ] + + +register_global(fp8x4_e4m3, Function(_ctor_template___nv_fp8x4_e4m3)) + + +def _from___nv_fp8x4_e4m3_to_float32x4_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK15__nv_fp8x4_e4m3cv6float4Ev_nbst(float4 &retval, __nv_fp8x4_e4m3 *self) { + retval = self->operator float4(); + return 0; + } + """ + + @lower_cast(_type_fp8x4_e4m3, float32x4) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK15__nv_fp8x4_e4m3cv6float4Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8x4_e4m3_to_float32x4_lower(shim_stream, shim_obj) + + +# Typing for fp8_e8m0 +class _type_class_fp8_e8m0(Number): + def __init__(self): + super().__init__(name="fp8_e8m0") + self.alignof_ = 1 + self.bitwidth = 1 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8_e8m0 = _type_class_fp8_e8m0() + + +# Make Python API for struct +fp8_e8m0 = type("fp8_e8m0", (), {"_nbtype": _type_fp8_e8m0}) + +as_numba_type.register(fp8_e8m0, _type_fp8_e8m0) + + +@register_model(_type_class_fp8_e8m0) +class _model_fp8_e8m0(PrimitiveModel): + def __init__(self, dmm, fe_type): + be_type = ir.IntType(fe_type.bitwidth) + super(_model_fp8_e8m0, self).__init__(dmm, fe_type, be_type) + + +def _lower__ZN13__nv_fp8_e8m0C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Ev_nbst(__nv_fp8_e8m0 *self ) { + new (self) __nv_fp8_e8m0(); + return 0; + } + """ + + @lower( + fp8_e8m0, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Ev(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1E6__half(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1E6__half_nbst(__nv_fp8_e8m0 *self , __half* f) { + new (self) __nv_fp8_e8m0(*f); + return 0; + } + """ + + @lower(fp8_e8m0, float16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1E6__half", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1E6__half(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1E13__nv_bfloat16(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1E13__nv_bfloat16_nbst(__nv_fp8_e8m0 *self , __nv_bfloat16* f) { + new (self) __nv_fp8_e8m0(*f); + return 0; + } + """ + + @lower(fp8_e8m0, bfloat16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1E13__nv_bfloat16", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1E13__nv_bfloat16(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Ef(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Ef_nbst(__nv_fp8_e8m0 *self , float* f) { + new (self) __nv_fp8_e8m0(*f); + return 0; + } + """ + + @lower(fp8_e8m0, float32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Ef", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Ef(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Ed(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Ed_nbst(__nv_fp8_e8m0 *self , double* f) { + new (self) __nv_fp8_e8m0(*f); + return 0; + } + """ + + @lower(fp8_e8m0, float64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Ed", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Ed(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Et(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Et_nbst(__nv_fp8_e8m0 *self , unsigned short* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, uint16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Et", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Et(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Ej(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Ej_nbst(__nv_fp8_e8m0 *self , unsigned int* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, uint32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Ej", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Ej(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Ey(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Ey_nbst(__nv_fp8_e8m0 *self , unsigned long long* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, uint64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Ey", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Ey(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Em(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Em_nbst(__nv_fp8_e8m0 *self , unsigned long* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, uint64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Em", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Em(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Es(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Es_nbst(__nv_fp8_e8m0 *self , short* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, int16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Es", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Es(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Ei(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Ei_nbst(__nv_fp8_e8m0 *self , int* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, int32) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Ei", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Ei(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1Ex(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1Ex_nbst(__nv_fp8_e8m0 *self , long long* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, int64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1Ex", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1Ex(shim_stream, shim_obj) + + +def _lower__ZN13__nv_fp8_e8m0C1El(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN13__nv_fp8_e8m0C1El_nbst(__nv_fp8_e8m0 *self , long* val) { + new (self) __nv_fp8_e8m0(*val); + return 0; + } + """ + + @lower(fp8_e8m0, int64) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN13__nv_fp8_e8m0C1El", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN13__nv_fp8_e8m0C1El(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8_e8m0(ConcreteTemplate): + key = globals()["fp8_e8m0"] + cases = [ + signature( + _type_fp8_e8m0, + ), + signature(_type_fp8_e8m0, float16), + signature(_type_fp8_e8m0, bfloat16), + signature(_type_fp8_e8m0, float32), + signature(_type_fp8_e8m0, float64), + signature(_type_fp8_e8m0, uint16), + signature(_type_fp8_e8m0, uint32), + signature(_type_fp8_e8m0, uint64), + signature(_type_fp8_e8m0, uint64), + signature(_type_fp8_e8m0, int16), + signature(_type_fp8_e8m0, int32), + signature(_type_fp8_e8m0, int64), + signature(_type_fp8_e8m0, int64), + ] + + +register_global(fp8_e8m0, Function(_ctor_template___nv_fp8_e8m0)) + + +def _from___nv_fp8_e8m0_to_float32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvfEv_nbst(float &retval, __nv_fp8_e8m0 *self) { + retval = self->operator float(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, float32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvfEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_float32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_bfloat16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cv13__nv_bfloat16Ev_nbst(__nv_bfloat16 &retval, __nv_fp8_e8m0 *self) { + retval = self->operator __nv_bfloat16(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, bfloat16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cv13__nv_bfloat16Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_bfloat16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_float64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvdEv_nbst(double &retval, __nv_fp8_e8m0 *self) { + retval = self->operator double(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, float64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvdEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_float64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_float16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cv6__halfEv_nbst(__half &retval, __nv_fp8_e8m0 *self) { + retval = self->operator __half(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, float16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cv6__halfEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_float16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_uint8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvhEv_nbst(unsigned char &retval, __nv_fp8_e8m0 *self) { + retval = self->operator unsigned char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, uint8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvhEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_uint8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_uint16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvtEv_nbst(unsigned short &retval, __nv_fp8_e8m0 *self) { + retval = self->operator unsigned short(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, uint16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvtEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_uint16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_uint32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvjEv_nbst(unsigned int &retval, __nv_fp8_e8m0 *self) { + retval = self->operator unsigned int(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, uint32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvjEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_uint32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_uint64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvmEv_nbst(unsigned long &retval, __nv_fp8_e8m0 *self) { + retval = self->operator unsigned long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, uint64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvmEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_uint64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_uint64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvyEv_nbst(unsigned long long &retval, __nv_fp8_e8m0 *self) { + retval = self->operator unsigned long long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, uint64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvyEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_uint64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_int8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvaEv_nbst(signed char &retval, __nv_fp8_e8m0 *self) { + retval = self->operator signed char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, int8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvaEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_int8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_int8_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvcEv_nbst(char &retval, __nv_fp8_e8m0 *self) { + retval = self->operator char(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, int8) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvcEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_int8_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_int16_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvsEv_nbst(short &retval, __nv_fp8_e8m0 *self) { + retval = self->operator short(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, int16) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvsEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_int16_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_int32_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cviEv_nbst(int &retval, __nv_fp8_e8m0 *self) { + retval = self->operator int(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, int32) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cviEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_int32_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_int64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvlEv_nbst(long &retval, __nv_fp8_e8m0 *self) { + retval = self->operator long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, int64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvlEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_int64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_int64_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvxEv_nbst(long long &retval, __nv_fp8_e8m0 *self) { + retval = self->operator long long(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, int64) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvxEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_int64_lower(shim_stream, shim_obj) + + +def _from___nv_fp8_e8m0_to_bool__lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK13__nv_fp8_e8m0cvbEv_nbst(bool &retval, __nv_fp8_e8m0 *self) { + retval = self->operator bool(); + return 0; + } + """ + + @lower_cast(_type_fp8_e8m0, bool_) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK13__nv_fp8_e8m0cvbEv", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8_e8m0_to_bool__lower(shim_stream, shim_obj) + + +# Typing for fp8x2_e8m0 +class _type_class_fp8x2_e8m0(Type): + def __init__(self): + super().__init__(name="fp8x2_e8m0") + self.alignof_ = 2 + self.bitwidth = 2 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8x2_e8m0 = _type_class_fp8x2_e8m0() + + +# Make Python API for struct +fp8x2_e8m0 = type("fp8x2_e8m0", (), {"_nbtype": _type_fp8x2_e8m0}) + +as_numba_type.register(fp8x2_e8m0, _type_fp8x2_e8m0) + + +@register_model(_type_class_fp8x2_e8m0) +class _model_fp8x2_e8m0(StructModel): + def __init__(self, dmm, fe_type): + members = [("__x", uint16)] + super().__init__(dmm, fe_type, members) + + +@register_attr +class _attr_typing_fp8x2_e8m0(AttributeTemplate): + key = _type_fp8x2_e8m0 + + def resolve___x(self, obj): + return uint16 + + +make_attribute_wrapper(_type_class_fp8x2_e8m0, "__x", "__x") + + +def _lower__ZN15__nv_fp8x2_e8m0C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e8m0C1Ev_nbst(__nv_fp8x2_e8m0 *self ) { + new (self) __nv_fp8x2_e8m0(); + return 0; + } + """ + + @lower( + fp8x2_e8m0, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e8m0C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e8m0C1Ev(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x2_e8m0C1E6float2(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e8m0C1E6float2_nbst(__nv_fp8x2_e8m0 *self , float2* f) { + new (self) __nv_fp8x2_e8m0(*f); + return 0; + } + """ + + @lower(fp8x2_e8m0, float32x2) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e8m0C1E6float2", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e8m0C1E6float2(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x2_e8m0C1E7double2(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x2_e8m0C1E7double2_nbst(__nv_fp8x2_e8m0 *self , double2* f) { + new (self) __nv_fp8x2_e8m0(*f); + return 0; + } + """ + + @lower(fp8x2_e8m0, float64x2) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x2_e8m0C1E7double2", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x2_e8m0C1E7double2(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8x2_e8m0(ConcreteTemplate): + key = globals()["fp8x2_e8m0"] + cases = [ + signature( + _type_fp8x2_e8m0, + ), + signature(_type_fp8x2_e8m0, float32x2), + signature(_type_fp8x2_e8m0, float64x2), + ] + + +register_global(fp8x2_e8m0, Function(_ctor_template___nv_fp8x2_e8m0)) + + +def _from___nv_fp8x2_e8m0_to_float32x2_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK15__nv_fp8x2_e8m0cv6float2Ev_nbst(float2 &retval, __nv_fp8x2_e8m0 *self) { + retval = self->operator float2(); + return 0; + } + """ + + @lower_cast(_type_fp8x2_e8m0, float32x2) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK15__nv_fp8x2_e8m0cv6float2Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8x2_e8m0_to_float32x2_lower(shim_stream, shim_obj) + + +# Typing for fp8x4_e8m0 +class _type_class_fp8x4_e8m0(Type): + def __init__(self): + super().__init__(name="fp8x4_e8m0") + self.alignof_ = 4 + self.bitwidth = 4 * 8 + + def can_convert_from(self, typingctx, other): + from numba.cuda.typeconv import Conversion + + if other in []: + return Conversion.safe + + +_type_fp8x4_e8m0 = _type_class_fp8x4_e8m0() + + +# Make Python API for struct +fp8x4_e8m0 = type("fp8x4_e8m0", (), {"_nbtype": _type_fp8x4_e8m0}) + +as_numba_type.register(fp8x4_e8m0, _type_fp8x4_e8m0) + + +@register_model(_type_class_fp8x4_e8m0) +class _model_fp8x4_e8m0(StructModel): + def __init__(self, dmm, fe_type): + members = [("__x", uint32)] + super().__init__(dmm, fe_type, members) + + +@register_attr +class _attr_typing_fp8x4_e8m0(AttributeTemplate): + key = _type_fp8x4_e8m0 + + def resolve___x(self, obj): + return uint32 + + +make_attribute_wrapper(_type_class_fp8x4_e8m0, "__x", "__x") + + +def _lower__ZN15__nv_fp8x4_e8m0C1Ev(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e8m0C1Ev_nbst(__nv_fp8x4_e8m0 *self ) { + new (self) __nv_fp8x4_e8m0(); + return 0; + } + """ + + @lower( + fp8x4_e8m0, + ) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e8m0C1Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e8m0C1Ev(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x4_e8m0C1E6float4(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e8m0C1E6float4_nbst(__nv_fp8x4_e8m0 *self , float4* f) { + new (self) __nv_fp8x4_e8m0(*f); + return 0; + } + """ + + @lower(fp8x4_e8m0, float32x4) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e8m0C1E6float4", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e8m0C1E6float4(shim_stream, shim_obj) + + +def _lower__ZN15__nv_fp8x4_e8m0C1E7double4(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN15__nv_fp8x4_e8m0C1E7double4_nbst(__nv_fp8x4_e8m0 *self , double4* f) { + new (self) __nv_fp8x4_e8m0(*f); + return 0; + } + """ + + @lower(fp8x4_e8m0, float64x4) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + ctor_callconv = FunctionCallConv( + itanium_mangled_name="_ZN15__nv_fp8x4_e8m0C1E7double4", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return ctor_callconv(builder, context, sig, args) + + +_lower__ZN15__nv_fp8x4_e8m0C1E7double4(shim_stream, shim_obj) + + +@register +class _ctor_template___nv_fp8x4_e8m0(ConcreteTemplate): + key = globals()["fp8x4_e8m0"] + cases = [ + signature( + _type_fp8x4_e8m0, + ), + signature(_type_fp8x4_e8m0, float32x4), + signature(_type_fp8x4_e8m0, float64x4), + ] + + +register_global(fp8x4_e8m0, Function(_ctor_template___nv_fp8x4_e8m0)) + + +def _from___nv_fp8x4_e8m0_to_float32x4_lower(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZNK15__nv_fp8x4_e8m0cv6float4Ev_nbst(float4 &retval, __nv_fp8x4_e8m0 *self) { + retval = self->operator float4(); + return 0; + } + """ + + @lower_cast(_type_fp8x4_e8m0, float32x4) + def impl(context, builder, fromty, toty, value): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZNK15__nv_fp8x4_e8m0cv6float4Ev", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + sig = signature(toty, fromty) + return callconv(builder, context, sig, [value]) + + +_from___nv_fp8x4_e8m0_to_float32x4_lower(shim_stream, shim_obj) + + +# Functions: + + +def cvt_double_to_fp8(): + pass + + +def _lower__ZL22__nv_cvt_double_to_fp8d17__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL22__nv_cvt_double_to_fp8d17__nv_saturation_t25__nv_fp8_interpretation_t_nbst(unsigned char &retval , double* x, __nv_saturation_t* saturate, __nv_fp8_interpretation_t* fp8_interpretation) { + retval = __nv_cvt_double_to_fp8(*x, *saturate, *fp8_interpretation); + return 0; + } + """ + + @lower( + cvt_double_to_fp8, + float64, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL22__nv_cvt_double_to_fp8d17__nv_saturation_t25__nv_fp8_interpretation_t", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL22__nv_cvt_double_to_fp8d17__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +) + + +def cvt_double2_to_fp8x2(): + pass + + +def _lower__ZL25__nv_cvt_double2_to_fp8x27double217__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL25__nv_cvt_double2_to_fp8x27double217__nv_saturation_t25__nv_fp8_interpretation_t_nbst(unsigned short &retval , double2* x, __nv_saturation_t* saturate, __nv_fp8_interpretation_t* fp8_interpretation) { + retval = __nv_cvt_double2_to_fp8x2(*x, *saturate, *fp8_interpretation); + return 0; + } + """ + + @lower( + cvt_double2_to_fp8x2, + float64x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL25__nv_cvt_double2_to_fp8x27double217__nv_saturation_t25__nv_fp8_interpretation_t", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL25__nv_cvt_double2_to_fp8x27double217__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +) + + +def cvt_float_to_fp8(): + pass + + +def _lower__ZL21__nv_cvt_float_to_fp8f17__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL21__nv_cvt_float_to_fp8f17__nv_saturation_t25__nv_fp8_interpretation_t_nbst(unsigned char &retval , float* x, __nv_saturation_t* saturate, __nv_fp8_interpretation_t* fp8_interpretation) { + retval = __nv_cvt_float_to_fp8(*x, *saturate, *fp8_interpretation); + return 0; + } + """ + + @lower( + cvt_float_to_fp8, + float32, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL21__nv_cvt_float_to_fp8f17__nv_saturation_t25__nv_fp8_interpretation_t", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL21__nv_cvt_float_to_fp8f17__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +) + + +def cvt_float2_to_fp8x2(): + pass + + +def _lower__ZL24__nv_cvt_float2_to_fp8x26float217__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL24__nv_cvt_float2_to_fp8x26float217__nv_saturation_t25__nv_fp8_interpretation_t_nbst(unsigned short &retval , float2* x, __nv_saturation_t* saturate, __nv_fp8_interpretation_t* fp8_interpretation) { + retval = __nv_cvt_float2_to_fp8x2(*x, *saturate, *fp8_interpretation); + return 0; + } + """ + + @lower( + cvt_float2_to_fp8x2, + float32x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL24__nv_cvt_float2_to_fp8x26float217__nv_saturation_t25__nv_fp8_interpretation_t", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL24__nv_cvt_float2_to_fp8x26float217__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +) + + +def cvt_bfloat16raw_to_fp8(): + pass + + +def _lower__ZL27__nv_cvt_bfloat16raw_to_fp817__nv_bfloat16_raw17__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL27__nv_cvt_bfloat16raw_to_fp817__nv_bfloat16_raw17__nv_saturation_t25__nv_fp8_interpretation_t_nbst(unsigned char &retval , __nv_bfloat16_raw* x, __nv_saturation_t* saturate, __nv_fp8_interpretation_t* fp8_interpretation) { + retval = __nv_cvt_bfloat16raw_to_fp8(*x, *saturate, *fp8_interpretation); + return 0; + } + """ + + @lower( + cvt_bfloat16raw_to_fp8, + bfloat16_raw_type, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL27__nv_cvt_bfloat16raw_to_fp817__nv_bfloat16_raw17__nv_saturation_t25__nv_fp8_interpretation_t", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL27__nv_cvt_bfloat16raw_to_fp817__nv_bfloat16_raw17__nv_saturation_t25__nv_fp8_interpretation_t_nbst( + shim_stream, shim_obj +) + + +def cvt_bfloat16raw_to_e8m0(): + pass + + +def _lower__ZL28__nv_cvt_bfloat16raw_to_e8m017__nv_bfloat16_raw17__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL28__nv_cvt_bfloat16raw_to_e8m017__nv_bfloat16_raw17__nv_saturation_t13cudaRoundMode_nbst(unsigned char &retval , __nv_bfloat16_raw* x, __nv_saturation_t* saturate, cudaRoundMode* rounding) { + retval = __nv_cvt_bfloat16raw_to_e8m0(*x, *saturate, *rounding); + return 0; + } + """ + + @lower( + cvt_bfloat16raw_to_e8m0, + bfloat16_raw_type, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL28__nv_cvt_bfloat16raw_to_e8m017__nv_bfloat16_raw17__nv_saturation_t13cudaRoundMode", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL28__nv_cvt_bfloat16raw_to_e8m017__nv_bfloat16_raw17__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +) + + +def cvt_float_to_e8m0(): + pass + + +def _lower__ZL22__nv_cvt_float_to_e8m0f17__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL22__nv_cvt_float_to_e8m0f17__nv_saturation_t13cudaRoundMode_nbst(unsigned char &retval , float* x, __nv_saturation_t* saturate, cudaRoundMode* rounding) { + retval = __nv_cvt_float_to_e8m0(*x, *saturate, *rounding); + return 0; + } + """ + + @lower( + cvt_float_to_e8m0, + float32, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL22__nv_cvt_float_to_e8m0f17__nv_saturation_t13cudaRoundMode", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL22__nv_cvt_float_to_e8m0f17__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +) + + +def cvt_float2_to_e8m0x2(): + pass + + +def _lower__ZL25__nv_cvt_float2_to_e8m0x26float217__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL25__nv_cvt_float2_to_e8m0x26float217__nv_saturation_t13cudaRoundMode_nbst(unsigned short &retval , float2* x, __nv_saturation_t* saturate, cudaRoundMode* rounding) { + retval = __nv_cvt_float2_to_e8m0x2(*x, *saturate, *rounding); + return 0; + } + """ + + @lower( + cvt_float2_to_e8m0x2, + float32x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL25__nv_cvt_float2_to_e8m0x26float217__nv_saturation_t13cudaRoundMode", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL25__nv_cvt_float2_to_e8m0x26float217__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +) + + +def cvt_double_to_e8m0(): + pass + + +def _lower__ZL23__nv_cvt_double_to_e8m0d17__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL23__nv_cvt_double_to_e8m0d17__nv_saturation_t13cudaRoundMode_nbst(unsigned char &retval , double* x, __nv_saturation_t* saturate, cudaRoundMode* rounding) { + retval = __nv_cvt_double_to_e8m0(*x, *saturate, *rounding); + return 0; + } + """ + + @lower( + cvt_double_to_e8m0, + float64, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL23__nv_cvt_double_to_e8m0d17__nv_saturation_t13cudaRoundMode", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL23__nv_cvt_double_to_e8m0d17__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +) + + +def cvt_double2_to_e8m0x2(): + pass + + +def _lower__ZL26__nv_cvt_double2_to_e8m0x27double217__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +): + shim_raw_str = """ + extern "C" __device__ int + _ZL26__nv_cvt_double2_to_e8m0x27double217__nv_saturation_t13cudaRoundMode_nbst(unsigned short &retval , double2* x, __nv_saturation_t* saturate, cudaRoundMode* rounding) { + retval = __nv_cvt_double2_to_e8m0x2(*x, *saturate, *rounding); + return 0; + } + """ + + @lower( + cvt_double2_to_e8m0x2, + float64x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL26__nv_cvt_double2_to_e8m0x27double217__nv_saturation_t13cudaRoundMode", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL26__nv_cvt_double2_to_e8m0x27double217__nv_saturation_t13cudaRoundMode_nbst( + shim_stream, shim_obj +) + + +def cvt_e8m0_to_bf16raw(): + pass + + +def _lower__ZL24__nv_cvt_e8m0_to_bf16rawh_nbst(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZL24__nv_cvt_e8m0_to_bf16rawh_nbst(__nv_bfloat16_raw &retval , unsigned char* x) { + retval = __nv_cvt_e8m0_to_bf16raw(*x); + return 0; + } + """ + + @lower(cvt_e8m0_to_bf16raw, uint8) + def impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + callconv = FunctionCallConv( + itanium_mangled_name="_ZL24__nv_cvt_e8m0_to_bf16rawh", + shim_writer=shim_writer, + shim_code=shim_raw_str, + ) + return callconv(builder, context, sig, args) + + +_lower__ZL24__nv_cvt_e8m0_to_bf16rawh_nbst(shim_stream, shim_obj) + + +@register +class _typing_cvt_double_to_fp8(ConcreteTemplate): + key = globals()["cvt_double_to_fp8"] + cases = [ + signature( + uint8, + float64, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + ] + + +register_global(cvt_double_to_fp8, types.Function(_typing_cvt_double_to_fp8)) + + +@register +class _typing_cvt_double2_to_fp8x2(ConcreteTemplate): + key = globals()["cvt_double2_to_fp8x2"] + cases = [ + signature( + uint16, + float64x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + ] + + +register_global( + cvt_double2_to_fp8x2, types.Function(_typing_cvt_double2_to_fp8x2) +) + + +@register +class _typing_cvt_float_to_fp8(ConcreteTemplate): + key = globals()["cvt_float_to_fp8"] + cases = [ + signature( + uint8, + float32, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + ] + + +register_global(cvt_float_to_fp8, types.Function(_typing_cvt_float_to_fp8)) + + +@register +class _typing_cvt_float2_to_fp8x2(ConcreteTemplate): + key = globals()["cvt_float2_to_fp8x2"] + cases = [ + signature( + uint16, + float32x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + ] + + +register_global( + cvt_float2_to_fp8x2, types.Function(_typing_cvt_float2_to_fp8x2) +) + + +@register +class _typing_cvt_bfloat16raw_to_fp8(ConcreteTemplate): + key = globals()["cvt_bfloat16raw_to_fp8"] + cases = [ + signature( + uint8, + bfloat16_raw_type, + IntEnumMember(saturation_t, int64), + IntEnumMember(fp8_interpretation_t, int64), + ) + ] + + +register_global( + cvt_bfloat16raw_to_fp8, types.Function(_typing_cvt_bfloat16raw_to_fp8) +) + + +@register +class _typing_cvt_bfloat16raw_to_e8m0(ConcreteTemplate): + key = globals()["cvt_bfloat16raw_to_e8m0"] + cases = [ + signature( + uint8, + bfloat16_raw_type, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + ] + + +register_global( + cvt_bfloat16raw_to_e8m0, types.Function(_typing_cvt_bfloat16raw_to_e8m0) +) + + +@register +class _typing_cvt_float_to_e8m0(ConcreteTemplate): + key = globals()["cvt_float_to_e8m0"] + cases = [ + signature( + uint8, + float32, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + ] + + +register_global(cvt_float_to_e8m0, types.Function(_typing_cvt_float_to_e8m0)) + + +@register +class _typing_cvt_float2_to_e8m0x2(ConcreteTemplate): + key = globals()["cvt_float2_to_e8m0x2"] + cases = [ + signature( + uint16, + float32x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + ] + + +register_global( + cvt_float2_to_e8m0x2, types.Function(_typing_cvt_float2_to_e8m0x2) +) + + +@register +class _typing_cvt_double_to_e8m0(ConcreteTemplate): + key = globals()["cvt_double_to_e8m0"] + cases = [ + signature( + uint8, + float64, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + ] + + +register_global(cvt_double_to_e8m0, types.Function(_typing_cvt_double_to_e8m0)) + + +@register +class _typing_cvt_double2_to_e8m0x2(ConcreteTemplate): + key = globals()["cvt_double2_to_e8m0x2"] + cases = [ + signature( + uint16, + float64x2, + IntEnumMember(saturation_t, int64), + IntEnumMember(cudaRoundMode, int64), + ) + ] + + +register_global( + cvt_double2_to_e8m0x2, types.Function(_typing_cvt_double2_to_e8m0x2) +) + + +@register +class _typing_cvt_e8m0_to_bf16raw(ConcreteTemplate): + key = globals()["cvt_e8m0_to_bf16raw"] + cases = [signature(bfloat16_raw_type, uint8)] + + +register_global( + cvt_e8m0_to_bf16raw, types.Function(_typing_cvt_e8m0_to_bf16raw) +) + + +# Aliases: + + +# Symbols: + + +_NBTYPE_SYMBOLS = [ + "_type_fp8_e5m2", + "_type_fp8x2_e5m2", + "_type_fp8x4_e5m2", + "_type_fp8_e4m3", + "_type_fp8x2_e4m3", + "_type_fp8x4_e4m3", + "_type_fp8_e8m0", + "_type_fp8x2_e8m0", + "_type_fp8x4_e8m0", +] + + +_RECORD_SYMBOLS = [ + "fp8_e5m2", + "fp8x2_e5m2", + "fp8x4_e5m2", + "fp8_e4m3", + "fp8x2_e4m3", + "fp8x4_e4m3", + "fp8_e8m0", + "fp8x2_e8m0", + "fp8x4_e8m0", +] + + +_FUNCTION_SYMBOLS = [ + "cvt_double_to_fp8", + "cvt_double2_to_fp8x2", + "cvt_float_to_fp8", + "cvt_float2_to_fp8x2", + "cvt_bfloat16raw_to_fp8", + "cvt_bfloat16raw_to_e8m0", + "cvt_float_to_e8m0", + "cvt_float2_to_e8m0x2", + "cvt_double_to_e8m0", + "cvt_double2_to_e8m0x2", + "cvt_e8m0_to_bf16raw", +] + + +_ENUM_SYMBOLS = ["saturation_t", "fp8_interpretation_t"] + +__all__ = _NBTYPE_SYMBOLS + _RECORD_SYMBOLS + _FUNCTION_SYMBOLS + _ENUM_SYMBOLS diff --git a/numba_cuda/numba/cuda/api.py b/numba_cuda/numba/cuda/api.py index 81c338806..31688d607 100644 --- a/numba_cuda/numba/cuda/api.py +++ b/numba_cuda/numba/cuda/api.py @@ -105,6 +105,14 @@ def is_bfloat16_supported(): return current_context().device.supports_bfloat16 +def is_fp8_supported(): + """Whether FP8 are supported. + + fp8 type and intrinsic functions are supported on devices with compute capability >= 8.9 + """ + return current_context().device.supports_fp8 + + @require_context def to_device(obj, stream=0, copy=True, to=None): """to_device(obj, stream=0, copy=True, to=None) @@ -538,6 +546,9 @@ def detect(): if os.name == "nt": attrs += [("Compute Mode", "TCC" if tcc else "WDDM")] attrs += [("FP32/FP64 Performance Ratio", fp32_to_fp64_ratio)] + attrs += [("FP16 Support", "Yes" if dev.supports_float16 else "No")] + attrs += [("BF16 Support", "Yes" if dev.supports_bfloat16 else "No")] + attrs += [("FP8 Support", "Yes" if dev.supports_fp8 else "No")] if cc < (3, 5): support = "[NOT SUPPORTED: CC < 3.5]" elif cc < (5, 0): diff --git a/numba_cuda/numba/cuda/cudadrv/driver.py b/numba_cuda/numba/cuda/cudadrv/driver.py index cc84a385e..f0c887e9c 100644 --- a/numba_cuda/numba/cuda/cudadrv/driver.py +++ b/numba_cuda/numba/cuda/cudadrv/driver.py @@ -619,6 +619,10 @@ def supports_float16(self): def supports_bfloat16(self): return self.compute_capability >= (8, 0) + @property + def supports_fp8(self): + return self.compute_capability >= (8, 9) + class BaseCUDAMemoryManager(object, metaclass=ABCMeta): """Abstract base class for External Memory Management (EMM) Plugins.""" diff --git a/numba_cuda/numba/cuda/fp8.py b/numba_cuda/numba/cuda/fp8.py new file mode 100644 index 000000000..4f27b08f9 --- /dev/null +++ b/numba_cuda/numba/cuda/fp8.py @@ -0,0 +1,12 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause +from numba.cuda._internal.cuda_fp8 import ( + typing_registry, + target_registry, +) + + +__all__ = [ + "typing_registry", + "target_registry", +] diff --git a/numba_cuda/numba/cuda/include/12/cuda_fp8.h b/numba_cuda/numba/cuda/include/12/cuda_fp8.h new file mode 100644 index 000000000..045f60b91 --- /dev/null +++ b/numba_cuda/numba/cuda/include/12/cuda_fp8.h @@ -0,0 +1,493 @@ +/* + * Copyright 2022-2024 NVIDIA Corporation. All rights reserved. + * + * NOTICE TO LICENSEE: + * + * This source code and/or documentation ("Licensed Deliverables") are + * subject to NVIDIA intellectual property rights under U.S. and + * international Copyright laws. + * + * These Licensed Deliverables contained herein is PROPRIETARY and + * CONFIDENTIAL to NVIDIA and is being provided under the terms and + * conditions of a form of NVIDIA software license agreement by and + * between NVIDIA and Licensee ("License Agreement") or electronically + * accepted by Licensee. Notwithstanding any terms or conditions to + * the contrary in the License Agreement, reproduction or disclosure + * of the Licensed Deliverables to any third party without the express + * written consent of NVIDIA is prohibited. + * + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE + * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS + * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND. + * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED + * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY, + * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY + * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY + * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS, + * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS + * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE + * OF THESE LICENSED DELIVERABLES. + * + * U.S. Government End Users. These Licensed Deliverables are a + * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT + * 1995), consisting of "commercial computer software" and "commercial + * computer software documentation" as such terms are used in 48 + * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government + * only as a commercial end item. Consistent with 48 C.F.R.12.212 and + * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all + * U.S. Government End Users acquire the Licensed Deliverables with + * only those rights set forth herein. + * + * Any use of the Licensed Deliverables in individual and commercial + * software must include, in the user documentation and internal + * comments to the code, the above Disclaimer and U.S. Government End + * Users Notice. + */ + +#ifndef __CUDA_FP8_H__ +#define __CUDA_FP8_H__ + +/* Set up function decorations */ +#if defined(__CUDACC__) +#define __CUDA_FP8_DECL__ static __device__ __inline__ +#define __CUDA_HOSTDEVICE_FP8__ __host__ __device__ +#define __CUDA_HOSTDEVICE_FP8_DECL__ static __host__ __device__ __inline__ +#else /* !defined(__CUDACC__) */ +#if defined(__GNUC__) +#define __CUDA_HOSTDEVICE_FP8_DECL__ static __attribute__((unused)) +#else +#define __CUDA_HOSTDEVICE_FP8_DECL__ static +#endif /* defined(__GNUC__) */ +#define __CUDA_HOSTDEVICE_FP8__ +#endif /* defined(__CUDACC_) */ + +#if !defined(_MSC_VER) && __cplusplus >= 201103L +#define __CPP_VERSION_AT_LEAST_11_FP8 +#elif _MSC_FULL_VER >= 190024210 && _MSVC_LANG >= 201103L +#define __CPP_VERSION_AT_LEAST_11_FP8 +#endif + +#if ( (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000) && (defined __CUDA_ARCH_FAMILY_SPECIFIC__) ) + #if (( \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ >= 1000) && \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ < 1100) \ + ) || ( \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ >= 1200) && \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ < 1300) \ + )) + #define __CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__ (1) + #else + // instructions unsupported, fallback to emulation + #define __CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__ (0) + #endif +#else /* ( (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000) && (defined __CUDA_ARCH_FAMILY_SPECIFIC__) ) */ + // instructions unsupported, fallback to emulation + #define __CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__ (0) +#endif /* ( (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000) && (defined __CUDA_ARCH_FAMILY_SPECIFIC__) ) */ + +// implicitly provided by NVRTC +#if !defined(__CUDACC_RTC__) +/* bring in enum cudaRoundMode */ +#include "device_types.h" +#endif /* !defined(__CUDACC_RTC__) */ + +/* bring in __half_raw data type */ +#include "cuda_fp16.h" +/* bring in __nv_bfloat16_raw data type */ +#include "cuda_bf16.h" + +// implicitly provided by NVRTC +#if !defined(__CUDACC_RTC__) +/* bring in float2, double4, etc vector types */ +#include "vector_types.h" +#endif /* !defined(__CUDACC_RTC__) */ + +/** + * \defgroup CUDA_MATH_INTRINSIC_FP8 FP8 Intrinsics + * This section describes fp8 intrinsic functions. + * To use these functions, include the header file \p cuda_fp8.h in your + * program. + * The following macros are available to help users selectively enable/disable + * various definitions present in the header file: + * - \p __CUDA_NO_FP8_CONVERSIONS__ - If defined, this macro will prevent any + * use of the C++ type conversions (converting constructors and conversion + * operators) defined in the header. + * - \p __CUDA_NO_FP8_CONVERSION_OPERATORS__ - If defined, this macro will + * prevent any use of the C++ conversion operators from \p fp8 to other types. + */ + +/** + * \defgroup CUDA_MATH_FP8_MISC FP8 Conversion and Data Movement + * \ingroup CUDA_MATH_INTRINSIC_FP8 + * To use these functions, include the header file \p cuda_fp8.h in your + * program. + */ + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief 8-bit \p unsigned \p integer + * type abstraction used for \p fp8 floating-point + * numbers storage. + */ +typedef unsigned char __nv_fp8_storage_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief 16-bit \p unsigned \p integer + * type abstraction used for storage of pairs of + * \p fp8 floating-point numbers. + */ +typedef unsigned short int __nv_fp8x2_storage_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief 32-bit \p unsigned \p integer + * type abstraction used for storage of tetrads of + * \p fp8 floating-point numbers. + */ +typedef unsigned int __nv_fp8x4_storage_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Enumerates the modes applicable when + * performing a narrowing conversion to \p fp8 destination types. + */ +typedef enum __nv_saturation_t { + /** + * Means no saturation to finite is performed when conversion + * results in rounding values outside the range of destination + * type. + * NOTE: for fp8 type of e4m3 kind, the results that are larger + * than the maximum representable finite number of the target + * format become NaN. + */ + __NV_NOSAT, + /** + * Means input larger than the maximum representable + * finite number MAXNORM of the target format round to the + * MAXNORM of the same sign as input. + */ + __NV_SATFINITE, +} __nv_saturation_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Enumerates the possible + * interpretations of the 8-bit values when referring to them as + * \p fp8 types. + */ +typedef enum __nv_fp8_interpretation_t { + __NV_E4M3, /**< Stands for \p fp8 numbers of \p e4m3 kind. */ + __NV_E5M2, /**< Stands for \p fp8 numbers of \p e5m2 kind. */ +} __nv_fp8_interpretation_t; + +/* Forward-declaration of C-style APIs */ + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p double precision \p x to \p fp8 type of the + * requested kind using round-to-nearest-even rounding and requested saturation + * mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_double_to_fp8(const double x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p double precision numbers packed + * in \p double2 \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_double2_to_fp8x2(const double2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p single precision \p x to \p fp8 type of the + * requested kind using round-to-nearest-even rounding and requested saturation + * mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_float_to_fp8(const float x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p single precision numbers packed + * in \p float2 \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_float2_to_fp8x2(const float2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p half precision \p x to \p fp8 type of the requested + * kind using round-to-nearest-even rounding and requested saturation mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_halfraw_to_fp8(const __half_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p half precision numbers packed + * in \p __half2_raw \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_halfraw2_to_fp8x2( + const __half2_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p nv_bfloat16 precision \p x to \p fp8 type of the + * requested kind using round-to-nearest-even rounding and requested saturation + * mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_fp8( + const __nv_bfloat16_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p nv_bfloat16 precision numbers packed + * in \p __nv_bfloat162_raw \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_bfloat16raw2_to_fp8x2( + const __nv_bfloat162_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p fp8 \p x of the specified kind + * to \p half precision. + * + * \details Converts input \p x of \p fp8 type of the kind specified by + * \p fp8_interpretation parameter + * to \p half precision. + * + * \returns + * - The \p __half_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __half_raw +__nv_cvt_fp8_to_halfraw(const __nv_fp8_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation); +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p fp8 values of the specified kind + * to a vector of two \p half precision values packed in \p __half2_raw + * structure. + * + * \details Converts input vector \p x of \p fp8 type of the kind specified by + * \p fp8_interpretation parameter + * to a vector of two \p half precision values and returns as \p __half2_raw + * structure. + * + * \returns + * - The \p __half2_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __half2_raw +__nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation); + + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p bfloat16 input into a scaling factor of \p e8m0 kind. + * + * \details Input number's absolute value is rounded to the closest power of two in the + * direction specified via \p rounding parameter. Rounded results that are + * smaller than the smallest representable target format number 2^-127 are then + * clipped to 2^-127. Results that are larger than the largest representable + * target format number 2^127 are either clipped to 2^127 if \p saturate equals + * to \p __NV_SATFINITE, or convert to \p NaN otherwise. \p NaN inputs convert + * into \p NaN output, encoded as \p 0xFF in the target format. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_e8m0(const __nv_bfloat16_raw x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts a pair of \p bfloat16 values into a pair of scaling factors of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_bfloat162raw_to_e8m0x2(const __nv_bfloat162_raw x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p float value into a scaling factor of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_float_to_e8m0(const float x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts a pair of \p float values into a pair of scaling factors of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_float2_to_e8m0x2(const float2 x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p double value into a scaling factor of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_double_to_e8m0(const double x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts a pair of \p double values into a pair of scaling factors of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_double2_to_e8m0x2(const double2 x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input scaling factor value of \p e8m0 kind into \p bfloat16. + * + * \details Input scales are exact powers of two or a \p NaN value, + * also representable in the target format. + * + * \returns + * - The \p __nv_bfloat16_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat16_raw __nv_cvt_e8m0_to_bf16raw(const __nv_fp8_storage_t x); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input pair of scaling factors of \p e8m0 kind into a pair of \p bfloat16 values. + * + * \returns + * - The \p __nv_bfloat162_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat162_raw __nv_cvt_e8m0x2_to_bf162raw(const __nv_fp8x2_storage_t x); + +#if defined(__cplusplus) + +#define __CUDA_FP8_TYPES_EXIST__ + +/* Forward-declaration of structures defined in "cuda_fp8.hpp" */ +struct __nv_fp8_e5m2; +struct __nv_fp8x2_e5m2; +struct __nv_fp8x4_e5m2; + +struct __nv_fp8_e4m3; +struct __nv_fp8x2_e4m3; +struct __nv_fp8x4_e4m3; + +struct __nv_fp8_e8m0; +struct __nv_fp8x2_e8m0; +struct __nv_fp8x4_e8m0; + +#endif /* defined(__cplusplus) */ + +#include "cuda_fp8.hpp" + +#undef __CUDA_FP8_DECL__ +#undef __CUDA_HOSTDEVICE_FP8__ +#undef __CUDA_HOSTDEVICE_FP8_DECL__ + +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) +#undef __CPP_VERSION_AT_LEAST_11_FP8 +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#endif /* end of include guard: __CUDA_FP8_H__ */ diff --git a/numba_cuda/numba/cuda/include/12/cuda_fp8.hpp b/numba_cuda/numba/cuda/include/12/cuda_fp8.hpp new file mode 100644 index 000000000..3903e8aa5 --- /dev/null +++ b/numba_cuda/numba/cuda/include/12/cuda_fp8.hpp @@ -0,0 +1,2710 @@ +/* + * Copyright 2022-2024 NVIDIA Corporation. All rights reserved. + * + * NOTICE TO LICENSEE: + * + * This source code and/or documentation ("Licensed Deliverables") are + * subject to NVIDIA intellectual property rights under U.S. and + * international Copyright laws. + * + * These Licensed Deliverables contained herein is PROPRIETARY and + * CONFIDENTIAL to NVIDIA and is being provided under the terms and + * conditions of a form of NVIDIA software license agreement by and + * between NVIDIA and Licensee ("License Agreement") or electronically + * accepted by Licensee. Notwithstanding any terms or conditions to + * the contrary in the License Agreement, reproduction or disclosure + * of the Licensed Deliverables to any third party without the express + * written consent of NVIDIA is prohibited. + * + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE + * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS + * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND. + * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED + * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY, + * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY + * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY + * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS, + * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS + * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE + * OF THESE LICENSED DELIVERABLES. + * + * U.S. Government End Users. These Licensed Deliverables are a + * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT + * 1995), consisting of "commercial computer software" and "commercial + * computer software documentation" as such terms are used in 48 + * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government + * only as a commercial end item. Consistent with 48 C.F.R.12.212 and + * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all + * U.S. Government End Users acquire the Licensed Deliverables with + * only those rights set forth herein. + * + * Any use of the Licensed Deliverables in individual and commercial + * software must include, in the user documentation and internal + * comments to the code, the above Disclaimer and U.S. Government End + * Users Notice. + */ + +#if !defined(__CUDA_FP8_HPP__) +#define __CUDA_FP8_HPP__ + +#if !defined(__CUDA_FP8_H__) +#error "Do not include this file directly. Instead, include cuda_fp8.h." +#endif + +/* C++ header for ::std::memcpy (used for type punning in host-side + * implementations). When compiling as a CUDA source file memcpy is provided + * implicitly. !defined(__CUDACC__) implies !defined(__CUDACC_RTC__). + */ +#if defined(__cplusplus) && !defined(__CUDACC__) +#include +#elif !defined(__cplusplus) && !defined(__CUDACC__) +#include +#endif /* defined(__cplusplus) && !defined(__CUDACC__) */ + +// implicitly provided by NVRTC +#if !defined(__CUDACC_RTC__) +#include +#endif /* !defined(__CUDACC_RTC__) */ + +#if !defined(IF_DEVICE_OR_CUDACC) +#if defined(__CUDACC__) + #define IF_DEVICE_OR_CUDACC(d, c, f) NV_IF_ELSE_TARGET(NV_IS_DEVICE, d, c) +#else + #define IF_DEVICE_OR_CUDACC(d, c, f) NV_IF_ELSE_TARGET(NV_IS_DEVICE, d, f) +#endif +#endif + +/* + * Bring in the standard assertions header to enforce the subset + * of rounding modes supported by the APIs defined here. + * NOTE: NVRTC defines its own assert + */ +#if !defined (__CUDACC_RTC__) +#include +#endif + +/* Set up structure-alignment attribute */ +#if !(defined __CUDA_ALIGN__) +#if defined(__CUDACC__) +#define __CUDA_ALIGN__(align) __align__(align) +#else +/* Define alignment macro based on compiler type (cannot assume C11 "_Alignas" + * is available) */ +#if __cplusplus >= 201103L +#define __CUDA_ALIGN__(n) \ + alignas(n) /* C++11 kindly gives us a keyword for this */ +#else /* !defined(__CPP_VERSION_AT_LEAST_11_FP8)*/ +#if defined(__GNUC__) +#define __CUDA_ALIGN__(n) __attribute__((aligned(n))) +#elif defined(_MSC_VER) +#define __CUDA_ALIGN__(n) __declspec(align(n)) +#else +#define __CUDA_ALIGN__(n) +#endif /* defined(__GNUC__) */ +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ +#endif /* defined(__CUDACC__) */ +#endif /* !(defined __CUDA_ALIGN__) */ + +#if !(defined __CPP_VERSION_AT_LEAST_11_FP8) +/* need c++11 for explicit operators */ +#define __CUDA_NO_FP8_CONVERSION_OPERATORS__ +#endif + +#if !(defined __DOXYGEN_ONLY__) + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_double_to_fp8(const double x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + unsigned char res; + unsigned long long int xbits; + +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&xbits, &x, sizeof(x)); +#else + (void)::std::memcpy(&xbits, &x, sizeof(x)); +#endif + unsigned char FP8_MAXNORM; + unsigned char FP8_MANTISSA_MASK; + unsigned short int FP8_EXP_BIAS; + unsigned long long int FP8_SIGNIFICAND_BITS; + const unsigned long long int DP_INF_BITS = 0x7FF0000000000000ULL; + unsigned long long int FP8_MINDENORM_O2; + unsigned long long int FP8_OVERFLOW_THRESHOLD; + unsigned long long int FP8_MINNORM; + + if (fp8_interpretation == __NV_E4M3) { + FP8_EXP_BIAS = 7U; + FP8_SIGNIFICAND_BITS = 4ULL; + FP8_MANTISSA_MASK = 0x7U; + FP8_MINDENORM_O2 = 0x3F50000000000000ULL; // mindenorm/2 = 2^-10 + FP8_OVERFLOW_THRESHOLD = + 0x407D000000000000ULL; // maxnorm + 1/2ulp = 0x1.Cp+8 + 0x1p+4 + FP8_MAXNORM = 0x7EU; + FP8_MINNORM = 0x3F90000000000000ULL; // minnorm = 2^-6 + } else { //__NV_E5M2 + FP8_EXP_BIAS = 15U; + FP8_SIGNIFICAND_BITS = 3ULL; + FP8_MANTISSA_MASK = 0x3U; + FP8_MINDENORM_O2 = 0x3EE0000000000000ULL; // mindenorm/2 = 2^-17 + FP8_OVERFLOW_THRESHOLD = + 0x40EE000000000000ULL - + 1ULL; // maxnorm + 1/2ulp = 0x1.Ep+15, and -1 to have common code + FP8_MAXNORM = 0x7BU; + FP8_MINNORM = 0x3F10000000000000ULL; // minnorm = 2^-14 + } + + // 1/2 LSB of the target format, positioned in double precision mantissa + // helpful in midpoints detection during round-to-nearest-even step + const unsigned long long int FP8_DP_HALF_ULP = + (unsigned long long int)1ULL << (53ULL - FP8_SIGNIFICAND_BITS - 1ULL); + // prepare sign bit in target format + unsigned char sign = (unsigned char)((xbits >> 63ULL) << 7U); + // prepare exponent field in target format + unsigned char exp = + (unsigned char)((((unsigned short int)(xbits >> 52ULL)) & 0x7FFU) - + 1023U + FP8_EXP_BIAS); + // round mantissa to target format width, rounding towards zero + unsigned char mantissa = + (unsigned char)(xbits >> (53ULL - FP8_SIGNIFICAND_BITS)) & + FP8_MANTISSA_MASK; + unsigned long long int absx = xbits & 0x7FFFFFFFFFFFFFFFULL; + + if (absx <= FP8_MINDENORM_O2) { + // zero or underflow + res = 0U; + } else if (absx > DP_INF_BITS) { + // NaN + if (fp8_interpretation == __NV_E4M3) { + res = 0x7FU; + } else { + // NaN --> QNaN + res = 0x7EU | mantissa; + } + } else if (absx > FP8_OVERFLOW_THRESHOLD) { + if (saturate == __NV_SATFINITE) { + res = FP8_MAXNORM; + } else { + // __NV_NOSAT + if (fp8_interpretation == __NV_E4M3) { + // no Inf in E4M3 + res = 0x7FU; // NaN + } else { + res = 0x7CU; // Inf in E5M2 + } + } + } else if (absx >= FP8_MINNORM) { + res = (unsigned char)((exp << (FP8_SIGNIFICAND_BITS - 1U)) | mantissa); + // rounded-off bits + unsigned long long int round = + xbits & ((FP8_DP_HALF_ULP << 1ULL) - 1ULL); + // round-to-nearest-even adjustment + if ((round > FP8_DP_HALF_ULP) || + ((round == FP8_DP_HALF_ULP) && (mantissa & 1U))) { + res = (unsigned char)(res + 1U); + } + } else // Denormal range + { + unsigned char shift = (unsigned char)(1U - exp); + // add implicit leading bit + mantissa |= (unsigned char)(1U << (FP8_SIGNIFICAND_BITS - 1U)); + // additional round-off due to denormalization + res = (unsigned char)(mantissa >> shift); + + // rounded-off bits, including implicit leading bit + unsigned long long int round = + (xbits | ((unsigned long long int)1ULL << (53ULL - 1ULL))) & + ((FP8_DP_HALF_ULP << (shift + 1ULL)) - 1ULL); + // round-to-nearest-even adjustment + if ((round > (FP8_DP_HALF_ULP << shift)) || + ((round == (FP8_DP_HALF_ULP << shift)) && (res & 1U))) { + res = (unsigned char)(res + 1U); + } + } + + res |= sign; + + return (__nv_fp8_storage_t)res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_double2_to_fp8x2(const double2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8x2_storage_t storage = (__nv_fp8x2_storage_t)__nv_cvt_double_to_fp8( + x.y, saturate, fp8_interpretation); + storage = (__nv_fp8x2_storage_t)(storage << 8U); + storage = (__nv_fp8x2_storage_t)(storage | + __nv_cvt_double_to_fp8( + x.x, saturate, fp8_interpretation)); + return storage; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_float_to_fp8(const float x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8_storage_t res = 0U; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + __nv_fp8x2_storage_t storage; + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x), "f"(0.0f)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x), "f"(0.0f)); + } + res = (__nv_fp8_storage_t)storage; + } else +#endif + { + unsigned int xbits; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&xbits, &x, sizeof(x)); +#else + (void)::std::memcpy(&xbits, &x, sizeof(x)); +#endif + + // isnan + if ((xbits & 0x7FFFFFFFU) > 0x7F800000U) { + // Canonical NaN + xbits = 0x7FFFFFFFU; + } + + float fx; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&fx, &xbits, sizeof(xbits)); +#else + (void)::std::memcpy(&fx, &xbits, sizeof(xbits)); +#endif + + const double dx = (double)fx; + res = __nv_cvt_double_to_fp8(dx, saturate, fp8_interpretation); + } + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_float2_to_fp8x2(const float2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8x2_storage_t storage; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x.x), "f"(x.y)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x.x), "f"(x.y)); + } + } else +#endif + { + storage = (__nv_fp8x2_storage_t)__nv_cvt_float_to_fp8( + x.y, saturate, fp8_interpretation); + storage = (__nv_fp8x2_storage_t)(storage << 8U); + storage = (__nv_fp8x2_storage_t)(storage | __nv_cvt_float_to_fp8( + x.x, saturate, + fp8_interpretation)); + } + return storage; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float +__internal_halfraw_to_float(const __half_raw x) { + float f; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 530) + asm("{cvt.f32.f16 %0, %1;}\n" : "=f"(f) : "h"(x.x)); +#else + const unsigned int ux = (unsigned int)x.x; + unsigned int sign = (ux >> 15U) & 1U; + unsigned int exponent = (ux >> 10U) & 0x1fU; + unsigned int mantissa = (ux & 0x3ffU) << 13U; + if (exponent == 0x1fU) { /* NaN or Inf */ + /* discard sign of a NaN */ + sign = ((mantissa != 0U) ? (sign >> 1U) : sign); + mantissa = ((mantissa != 0U) ? 0x7fffffU : 0U); + exponent = 0xffU; + } else if (exponent == 0U) { /* Denorm or Zero */ + if (mantissa != 0U) { + unsigned int msb; + exponent = 0x71U; + do { + msb = (mantissa & 0x400000U); + mantissa <<= 1U; /* normalize */ + --exponent; + } while (msb == 0U); + mantissa &= 0x7fffffU; /* 1.mantissa is implicit */ + } + } else { + exponent += 0x70U; + } + const unsigned int u = ((sign << 31U) | (exponent << 23U) | mantissa); +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&f, &u, sizeof(u)); +#else + (void)::std::memcpy(&f, &u, sizeof(u)); +#endif +#endif /* (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 530) */ + return f; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float2 +__internal_halfraw2_to_float2(const __half2_raw x) { + __half_raw raw; + float2 res; + raw.x = x.x; + res.x = __internal_halfraw_to_float(raw); + raw.x = x.y; + res.y = __internal_halfraw_to_float(raw); + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_halfraw_to_fp8(const __half_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8_storage_t res = 0U; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + unsigned int half2_storage = (unsigned int)(x.x); + __nv_fp8x2_storage_t tmp; + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } + res = (__nv_fp8_storage_t)tmp; + } else +#endif + { + float fx = __internal_halfraw_to_float(x); + res = __nv_cvt_float_to_fp8(fx, saturate, fp8_interpretation); + } + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_halfraw2_to_fp8x2( + const __half2_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8x2_storage_t tmp; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + unsigned int half2_storage; + (void)memcpy(&half2_storage, &x, sizeof(x)); + + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } + } else +#endif + { + __half_raw raw; + raw.x = x.x; + __nv_fp8_storage_t lo = + __nv_cvt_halfraw_to_fp8(raw, saturate, fp8_interpretation); + raw.x = x.y; + __nv_fp8_storage_t hi = + __nv_cvt_halfraw_to_fp8(raw, saturate, fp8_interpretation); + tmp = hi; + tmp = (__nv_fp8x2_storage_t)(tmp << 8U); + tmp = (__nv_fp8x2_storage_t)(tmp | lo); + } + return tmp; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float +__internal_bf16raw_to_float(const __nv_bfloat16_raw x) { + const unsigned int ux = ((unsigned int)x.x) << 16U; + float fx; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&fx, &ux, sizeof(ux)); +#else + (void)::std::memcpy(&fx, &ux, sizeof(ux)); +#endif + return fx; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_fp8( + const __nv_bfloat16_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + const float fx = __internal_bf16raw_to_float(x); + const __nv_fp8_storage_t res = + __nv_cvt_float_to_fp8(fx, saturate, fp8_interpretation); + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_bfloat16raw2_to_fp8x2( + const __nv_bfloat162_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_bfloat16_raw raw; + raw.x = x.y; + __nv_fp8x2_storage_t storage = + (__nv_fp8x2_storage_t)__nv_cvt_bfloat16raw_to_fp8(raw, saturate, + fp8_interpretation); + storage = (__nv_fp8x2_storage_t)(storage << 8U); + raw.x = x.x; + storage = (__nv_fp8x2_storage_t)(storage | + __nv_cvt_bfloat16raw_to_fp8( + raw, saturate, fp8_interpretation)); + return storage; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __half2_raw +__nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation); + +__CUDA_HOSTDEVICE_FP8_DECL__ __half_raw +__nv_cvt_fp8_to_halfraw(const __nv_fp8_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation) { + __half_raw res; + res.x = 0U; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + res.x = + __nv_cvt_fp8x2_to_halfraw2((__nv_fp8x2_storage_t)x, fp8_interpretation) + .x; +#else + unsigned short int ur = (unsigned short int)x; + ur = (unsigned short int)(ur << 8U); + + if (fp8_interpretation == __NV_E5M2) { + if ((ur & 0x7FFFU) > 0x7C00U) { + /* If NaN, return canonical NaN */ + ur = 0x7FFFU; + } + } else { // __NV_E4M3 + unsigned short int sign = ur & 0x8000U; + unsigned short int exponent = + (unsigned short int)(((ur & 0x7800U) >> 1U) + 0x2000U); + unsigned short int mantissa = (ur & 0x0700U) >> 1U; + unsigned char absx = 0x7FU & (unsigned char)x; + + if (absx == 0x7FU) // NaN + { + ur = 0x7FFFU; // fp16 canonical NaN, discard sign + } else if (exponent == 0x2000U) { + // zero or denormal + if (mantissa != 0U) { + // normalize + mantissa = (unsigned short int)(mantissa << 1U); + while ((mantissa & 0x0400U) == 0U) { + mantissa = (unsigned short int)(mantissa << 1U); + exponent = (unsigned short int)(exponent - 0x0400U); + } + // discard implicit leading bit + mantissa &= 0x03FFU; + } else { // Zero + exponent = 0U; + } + + ur = (sign | exponent) | mantissa; + } else { + ur = (sign | exponent) | mantissa; + } + } + res.x = ur; +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __half2_raw +__nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation) { + __half2_raw res; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + unsigned int half2_storage; + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.f16x2.e5m2x2 %0, %1;}\n" : "=r"(half2_storage) : "h"(x)); + } else { + asm("{cvt.rn.f16x2.e4m3x2 %0, %1;}\n" : "=r"(half2_storage) : "h"(x)); + } + (void)memcpy(&res, &half2_storage, sizeof(half2_storage)); +#else + res.x = + __nv_cvt_fp8_to_halfraw((__nv_fp8_storage_t)x, fp8_interpretation).x; + res.y = __nv_cvt_fp8_to_halfraw((__nv_fp8_storage_t)(x >> 8U), + fp8_interpretation) + .x; +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8_storage_t __nv_cvt_bfloat16raw_to_e8m0(const __nv_bfloat16_raw x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short ures = 0U; + unsigned in = (unsigned)(x.x); + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } + return (__nv_fp8_storage_t)ures; +#else + // extract exponent bits, provides non-saturated result in RZ + __nv_fp8_storage_t res = (unsigned char)(x.x >> 7U); + + if (rounding == cudaRoundPosInf) { + // round-up if mantissa non-zero and |x| > 2^-127 and finite + if ((x.x & 0x007FU) && ((x.x & 0x7FFFU) > 0x0040U) && ((x.x & 0x7FFFU) < 0x7F80U)) res++; + } + + // Handle saturation of non-NaN large inputs to finite + if (saturate == __NV_SATFINITE) { + // non-NaN, Overflow --> Max + if (((x.x & 0x7FFFU) <= 0x7F80U) && (res == 0xFFU)) + { + res--; + } + } + return res; +#endif +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8x2_storage_t __nv_cvt_bfloat162raw_to_e8m0x2(const __nv_bfloat162_raw x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + assert((rounding == cudaRoundZero) || (rounding == cudaRoundPosInf)); +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + __nv_fp8x2_storage_t ures = 0U; + unsigned in = (unsigned)(x.x) | ((unsigned)(x.y) << (unsigned)16U); + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } + return ures; +#else + __nv_bfloat16_raw lo, hi; + lo.x = x.x; + hi.x = x.y; + __nv_fp8x2_storage_t ures = __nv_cvt_bfloat16raw_to_e8m0(hi, saturate, rounding); + ures <<= (unsigned short)8U; + ures |= __nv_cvt_bfloat16raw_to_e8m0(lo, saturate, rounding); + return ures; +#endif +} + +__CUDA_HOSTDEVICE_FP8_DECL__ unsigned int __internal_fp8_float_as_uint(const float f) +{ + unsigned int u; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&u, &f, sizeof(f)); +#else + (void)::std::memcpy(&u, &f, sizeof(f)); +#endif + return u; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float __internal_fp8_uint_as_float(const unsigned int u) +{ + float f; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&f, &u, sizeof(u)); +#else + (void)::std::memcpy(&f, &u, sizeof(u)); +#endif + return f; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat16_raw +__internal_float_to_bf16raw_rz(const float x) { + __nv_bfloat16_raw r; +NV_IF_ELSE_TARGET(NV_PROVIDES_SM_80, + asm("{ cvt.rz.bf16.f32 %0, %1;}\n" : "=h"(r.x) : "f"(x)); +, + unsigned int ux = __internal_fp8_float_as_uint(x); + if ((ux & 0x7FFFFFFFU) > 0x7f800000U) + { + // NaN + r.x = (unsigned short int)0x7FFFU; + } + else + { + r.x = (unsigned short int)(ux >> 16U); + } +) + return r; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat16_raw +__internal_float_to_bf16raw_ru(const float x) { + __nv_bfloat16_raw r; +NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90, + asm("{ cvt.rp.bf16.f32 %0, %1;}\n" : "=h"(r.x) : "f"(x)); +, + unsigned int ux = __internal_fp8_float_as_uint(x); + if ((ux & 0x7FFFFFFFU) > 0x7f800000U) + { + // NaN + r.x = (unsigned short int)0x7FFFU; + } + else if ((ux < 0x7f800000U) && ((ux & 0x0000FFFFU) != 0)) + { + // 0 <= x < +inf, round-up + r.x = (unsigned short int)((ux >> 16U) + 1U); + } + else { + // truncate others + r.x = (unsigned short int)(ux >> 16U); + } +) + return r; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8_storage_t __nv_cvt_float_to_e8m0(const float x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + assert((rounding == cudaRoundZero) || (rounding == cudaRoundPosInf)); + __nv_fp8_storage_t res = 0U; +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short ures = 0U; + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } + res = (__nv_fp8_storage_t)ures; +#else + if (rounding == cudaRoundZero) + { + res = __nv_cvt_bfloat16raw_to_e8m0(__internal_float_to_bf16raw_rz(x), saturate, rounding); + } + else + { //cudaRoundPosInf + float absx = __internal_fp8_uint_as_float((__internal_fp8_float_as_uint(x) << (unsigned)1U) >> (unsigned)1U); + res = __nv_cvt_bfloat16raw_to_e8m0(__internal_float_to_bf16raw_ru(absx), saturate, rounding); + } +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8x2_storage_t __nv_cvt_float2_to_e8m0x2(const float2 x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + assert((rounding == cudaRoundZero) || (rounding == cudaRoundPosInf)); + __nv_fp8x2_storage_t res = 0U; +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } +#else + res = __nv_cvt_float_to_e8m0(x.y, saturate, rounding); + res <<= (unsigned short)8U; + res |= __nv_cvt_float_to_e8m0(x.x, saturate, rounding); +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +float __internal_double_to_float_with_sticky(double x) +{ +#if (defined __CUDA_ARCH__) + // protect from ftz in device code + float f; + double d; + asm("{ cvt.rn.f32.f64 %0, %1;}\n" : "=f"(f) : "d"(x)); + asm("{ cvt.f64.f32 %0, %1;}\n" : "=d"(d) : "f"(f)); +#else + const float f = (float)x; + const double d = (double)f; +#endif + unsigned int u = __internal_fp8_float_as_uint(f); + int x_is_not_nan = ((u << (unsigned)1U) <= (unsigned)0xFF000000U) ? 1 : 0; + + if ((x > 0.0) && (d > x)) { + u--; + } + if ((x < 0.0) && (d < x)) { + u--; + } + if ((d != x) && (x_is_not_nan == 1)) { + u |= 1U; + } + return __internal_fp8_uint_as_float(u); +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8_storage_t __nv_cvt_double_to_e8m0(const double x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + float fx_with_sticky = __internal_double_to_float_with_sticky(x); + __nv_fp8_storage_t res = __nv_cvt_float_to_e8m0(fx_with_sticky, saturate, rounding); + + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8x2_storage_t __nv_cvt_double2_to_e8m0x2(const double2 x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + float2 f; + f.x = __internal_double_to_float_with_sticky(x.x); + f.y = __internal_double_to_float_with_sticky(x.y); + return __nv_cvt_float2_to_e8m0x2(f, saturate, rounding); +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +unsigned short __internal_e8m0_to_bf16(const __nv_fp8_storage_t x) +{ + unsigned short res; + // shift bias exponent bits into place + res = ((unsigned short)x) << 7U; + + if (x == 0xFFU) { + res = 0x7FFFU; // NaN --> Canonical QNaN + } else if (x == 0U) { + res = 0x0040U; // 2^-127 + } + + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_bfloat16_raw __nv_cvt_e8m0_to_bf16raw(const __nv_fp8_storage_t x) +{ + __nv_bfloat16_raw res; + +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short in = (unsigned short)x; + unsigned hr = 0U; + asm("{cvt.rn.bf16x2.ue8m0x2 %0, %1;}\n" + : "=r"(hr) + : "h"(in)); + + res.x = (unsigned short)hr; +#else + res.x = __internal_e8m0_to_bf16(x); +#endif + + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_bfloat162_raw __nv_cvt_e8m0x2_to_bf162raw(const __nv_fp8x2_storage_t x) +{ + __nv_bfloat162_raw res; + +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short in = (unsigned short)x; + unsigned hr = 0U; + asm("{cvt.rn.bf16x2.ue8m0x2 %0, %1;}\n" + : "=r"(hr) + : "h"(in)); + + res.x = (unsigned short)hr; + res.y = (unsigned short)(hr >> (unsigned)16U); +#else + + res.x = __internal_e8m0_to_bf16((__nv_fp8_storage_t)x); + res.y = __internal_e8m0_to_bf16((__nv_fp8_storage_t)(x >> (unsigned short)8U)); + +#endif + + return res; +} + +#endif /* !(defined __DOXYGEN_ONLY__) */ + +/* All other definitions in this file are only visible to C++ compilers */ +#if defined(__cplusplus) + +/** + * \defgroup CUDA_MATH_FP8_E5M2_STRUCT C++ struct for handling fp8 data type of e5m2 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8_E5M2_STRUCT + * \brief __nv_fp8_e5m2 datatype + * + * \details This structure implements the datatype for handling + * \p fp8 floating-point numbers of \p e5m2 kind: + * with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(1) __nv_fp8_e5m2 { + public: + /** + * \ingroup CUDA_MATH_FP8_E5M2_STRUCT + * Storage variable contains the \p fp8 floating-point data. + */ + __nv_fp8_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8_e5m2() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider FP types */ + /* Note we do avoid constructor init-list because of special host/device + * compilation rules */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const __half f) { + __x = __nv_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat16 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const __nv_bfloat16 f) { + __x = __nv_cvt_bfloat16raw_to_fp8(static_cast<__nv_bfloat16_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const float f) { + __x = __nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const double f) { + __x = __nv_cvt_double_to_fp8(f, __NV_SATFINITE, __NV_E5M2); + } + + /* Converts from integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p short \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e5m2(const unsigned short int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const unsigned int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const unsigned long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p long \p int data type, relies on + * \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e5m2(const unsigned long long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p short \p int data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const short int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p int data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p int data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const long long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening FP converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half() const { + return static_cast<__half>(__nv_cvt_fp8_to_halfraw(__x, __NV_E5M2)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float() const { + return __internal_halfraw_to_float( + __nv_cvt_fp8_to_halfraw(__x, __NV_E5M2)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat16 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat16() const { + return __float2bfloat16_rz(float(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p double data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator double() const { + return static_cast(float(*this)); + } + + /* Convert to integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p char data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned char() const { + unsigned char i; + const float f = float(*this); + const unsigned char max_val = 0xFFU; + const unsigned char min_val = 0U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) > 0x7CU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p short \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned short int() const { + return __half2ushort_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned int() const { + return __half2uint_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long int() const { + unsigned long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(unsigned long) == sizeof(unsigned long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ull_rz(__half(*this))); + } + else + { + retval = static_cast(__half2uint_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p long \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000ULL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long long int() const { + return __half2ull_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p signed \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator signed char() const { + signed char i; + const float f = float(*this); + const signed char max_val = (signed char)0x7FU; + const signed char min_val = (signed char)0x80U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) > 0x7CU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to an implementation defined \p char data type. + * + * Detects signedness of the \p char type and proceeds accordingly, see + * further details in signed and unsigned char operators. + + * Clamps inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator char() const { + char value; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (((char)-1) < (char)0) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + value = static_cast(static_cast(*this)); + } + else + { + value = static_cast(static_cast(*this)); + } + return value; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p short \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator short int() const { + return __half2short_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator int() const { + return __half2int_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long int() const { + long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(long) == sizeof(long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ll_rz(__half(*this))); + } + else + { + retval = static_cast(__half2int_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000LL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long long int() const { + return __half2ll_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p bool data type. + * +0 and -0 inputs convert to \p false. + * Non-zero inputs convert to \p true. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator bool() const { + return (__x & 0x7FU) != 0U; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X2_E5M2_STRUCT C++ struct for handling vector type of two fp8 values of e5m2 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X2_E5M2_STRUCT + * \brief __nv_fp8x2_e5m2 datatype + * + * \details This structure implements the datatype for handling two + * \p fp8 floating-point numbers of \p e5m2 kind each: + * with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(2) __nv_fp8x2_e5m2 { + public: + /** + * \ingroup CUDA_MATH_FP8X2_E5M2_STRUCT + * Storage variable contains the vector of two \p fp8 floating-point data + * values. + */ + __nv_fp8x2_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x2_e5m2() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const __half2 f) { + __x = __nv_cvt_halfraw2_to_fp8x2(static_cast<__half2_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat162 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const __nv_bfloat162 f) { + __x = __nv_cvt_bfloat16raw2_to_fp8x2(static_cast<__nv_bfloat162_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const float2 f) { + __x = __nv_cvt_float2_to_fp8x2(f, __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const double2 f) { + __x = __nv_cvt_double2_to_fp8x2(f, __NV_SATFINITE, __NV_E5M2); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half2() const { + return static_cast<__half2>(__nv_cvt_fp8x2_to_halfraw2(__x, __NV_E5M2)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float2() const { + return __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(__x, __NV_E5M2)); + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +__CUDA_HOSTDEVICE_FP8_DECL__ unsigned int +__internal_pack_u16x2_to_u32(const unsigned short int src_lo, + const unsigned short int src_hi) { + unsigned int dst; +#if (defined __CUDACC__) && (defined __CUDA_ARCH__) + asm("{ mov.b32 %0, {%1,%2};}\n" : "=r"(dst) : "h"(src_lo), "h"(src_hi)); +#else + dst = (static_cast(src_hi) << 16U) | + static_cast(src_lo); +#endif + return dst; +} + +/** + * \defgroup CUDA_MATH_FP8X4_E5M2_STRUCT C++ struct for handling vector type of four fp8 values of e5m2 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X4_E5M2_STRUCT + * \brief __nv_fp8x4_e5m2 datatype + * + * \details This structure implements the datatype for handling four + * \p fp8 floating-point numbers of \p e5m2 kind each: + * with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(4) __nv_fp8x4_e5m2 { + public: + /** + * \ingroup CUDA_MATH_FP8X4_E5M2_STRUCT + * Storage variable contains the vector of four \p fp8 floating-point data + * values. + */ + __nv_fp8x4_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x4_e5m2() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __half2 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const __half2 flo, + const __half2 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(flo), __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(fhi), __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __nv_bfloat162 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const __nv_bfloat162 flo, + const __nv_bfloat162 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(flo), __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(fhi), __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const float4 f) { + const float2 flo = {f.x, f.y}; + const float2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_float2_to_fp8x2(flo, __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_float2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const double4 f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float4 vector data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float4() const { + const __nv_fp8x2_storage_t slo = static_cast<__nv_fp8x2_storage_t>(__x); + const __nv_fp8x2_storage_t shi = + static_cast<__nv_fp8x2_storage_t>(__x >> 16U); + float2 rlo = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(slo, __NV_E5M2)); + float2 rhi = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(shi, __NV_E5M2)); + float4 res = {rlo.x, rlo.y, rhi.x, rhi.y}; + return res; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8_E4M3_STRUCT C++ struct for handling fp8 data type of e4m3 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8_E4M3_STRUCT + * \brief __nv_fp8_e4m3 datatype + * + * \details This structure implements the datatype for storing + * \p fp8 floating-point numbers of \p e4m3 kind: + * with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. + * The encoding doesn't support Infinity. + * NaNs are limited to 0x7F and 0xFF values. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(1) __nv_fp8_e4m3 { + public: + /** + * \ingroup CUDA_MATH_FP8_E4M3_STRUCT + * Storage variable contains the \p fp8 floating-point data. + */ + __nv_fp8_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8_e4m3() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider FP types */ + /* Note we do avoid constructor init-list because of special host/device + * compilation rules */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const __half f) { + __x = __nv_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat16 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const __nv_bfloat16 f) { + __x = __nv_cvt_bfloat16raw_to_fp8(static_cast<__nv_bfloat16_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const float f) { + __x = __nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const double f) { + __x = __nv_cvt_double_to_fp8(f, __NV_SATFINITE, __NV_E4M3); + } + + /* Converts from integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p short \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e4m3(const unsigned short int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const unsigned int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const unsigned long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p long \p int data type, relies on + * \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e4m3(const unsigned long long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p short \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const short int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p int data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const long long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening FP converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half() const { + return static_cast<__half>(__nv_cvt_fp8_to_halfraw(__x, __NV_E4M3)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float() const { + return __internal_halfraw_to_float( + __nv_cvt_fp8_to_halfraw(__x, __NV_E4M3)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat16 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat16() const { + return __float2bfloat16_rz(float(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p double data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator double() const { + return static_cast(float(*this)); + } + + /* Convert to integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p char data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned char() const { + unsigned char i; + const float f = float(*this); + const unsigned char max_val = 0xFFU; + const unsigned char min_val = 0U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) == 0x7FU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p short \p int data type. + * Clamps negative inputs to zero. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned short int() const { + return __half2ushort_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p int data type. + * Clamps negative inputs to zero. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned int() const { + return __half2uint_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long int() const { + unsigned long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(unsigned long) == sizeof(unsigned long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ull_rz(__half(*this))); + } + else + { + retval = static_cast(__half2uint_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p long \p int data type. + * Clamps negative inputs to zero. + * \p NaN inputs convert to \p 0x8000000000000000ULL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long long int() const { + return __half2ull_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p signed \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator signed char() const { + signed char i; + const float f = float(*this); + const signed char max_val = (signed char)0x7FU; + const signed char min_val = (signed char)0x80U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) == 0x7FU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to an implementation defined \p char data type. + * + * Detects signedness of the \p char type and proceeds accordingly, see + * further details in signed and unsigned char operators. + + * Clamps inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator char() const { + char value; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (((char)-1) < (char)0) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + value = static_cast(static_cast(*this)); + } + else + { + value = static_cast(static_cast(*this)); + } + return value; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p short \p int data type. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator short int() const { + return __half2short_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p int data type. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator int() const { + return __half2int_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long int() const { + long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(long) == sizeof(long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ll_rz(__half(*this))); + } + else + { + retval = static_cast(__half2int_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p long \p int data type. + * \p NaN inputs convert to \p 0x8000000000000000LL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long long int() const { + return __half2ll_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p bool data type. + * +0 and -0 inputs convert to \p false. + * Non-zero inputs convert to \p true. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator bool() const { + return (__x & 0x7FU) != 0U; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X2_E4M3_STRUCT C++ struct for handling vector type of two fp8 values of e4m3 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X2_E4M3_STRUCT + * \brief __nv_fp8x2_e4m3 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of two \p fp8 values of \p e4m3 kind each: + * with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. + * The encoding doesn't support Infinity. + * NaNs are limited to 0x7F and 0xFF values. + */ +struct __CUDA_ALIGN__(2) __nv_fp8x2_e4m3 { + public: + /** + * \ingroup CUDA_MATH_FP8X2_E4M3_STRUCT + * Storage variable contains the vector of two \p fp8 floating-point data + * values. + */ + __nv_fp8x2_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x2_e4m3() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const __half2 f) { + __x = __nv_cvt_halfraw2_to_fp8x2(static_cast<__half2_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat162 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const __nv_bfloat162 f) { + __x = __nv_cvt_bfloat16raw2_to_fp8x2(static_cast<__nv_bfloat162_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const float2 f) { + __x = __nv_cvt_float2_to_fp8x2(f, __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const double2 f) { + __x = __nv_cvt_double2_to_fp8x2(f, __NV_SATFINITE, __NV_E4M3); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half2() const { + return static_cast<__half2>(__nv_cvt_fp8x2_to_halfraw2(__x, __NV_E4M3)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float2() const { + return __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(__x, __NV_E4M3)); + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X4_E4M3_STRUCT C++ struct for handling vector type of four fp8 values of e4m3 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X4_E4M3_STRUCT + * \brief __nv_fp8x4_e4m3 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of four \p fp8 values of \p e4m3 kind each: + * with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. + * The encoding doesn't support Infinity. + * NaNs are limited to 0x7F and 0xFF values. + */ +struct __CUDA_ALIGN__(4) __nv_fp8x4_e4m3 { + public: + /** + * \ingroup CUDA_MATH_FP8X4_E4M3_STRUCT + * Storage variable contains the vector of four \p fp8 floating-point data + * values. + */ + __nv_fp8x4_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x4_e4m3() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __half2 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const __half2 flo, + const __half2 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(flo), __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(fhi), __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __nv_bfloat162 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const __nv_bfloat162 flo, + const __nv_bfloat162 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(flo), __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(fhi), __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const float4 f) { + const float2 flo = {f.x, f.y}; + const float2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_float2_to_fp8x2(flo, __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_float2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const double4 f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float4 vector data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float4() const { + const __nv_fp8x2_storage_t slo = static_cast<__nv_fp8x2_storage_t>(__x); + const __nv_fp8x2_storage_t shi = + static_cast<__nv_fp8x2_storage_t>(__x >> 16U); + float2 rlo = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(slo, __NV_E4M3)); + float2 rhi = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(shi, __NV_E4M3)); + float4 res = {rlo.x, rlo.y, rhi.x, rhi.y}; + return res; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \ingroup CUDA_MATH_FP8_E8M0_STRUCT + * \brief __nv_fp8_e8m0 datatype + * + * \details This structure implements the datatype for handling + * 8-bit scale factors of \p e8m0 kind: interpreted as powers of two + * with biased exponent. Bias equals to 127, so numbers 0 through 254 + * represent 2^-127 through 2^127. Number \p 0xFF = 255 is reserved + * for NaN. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(1) __nv_fp8_e8m0 { + public: + /** + * \ingroup CUDA_MATH_FP8_E8M0_STRUCT + * Storage variable contains the 8-bit scale data. + */ + __nv_fp8_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8_e8m0() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider FP types */ + /* Note we do avoid constructor init-list because of special host/device + * compilation rules */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half data type, relies on \p __NV_SATFINITE + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_float_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const __half f) { + __x = __nv_cvt_float_to_e8m0(__internal_halfraw_to_float(static_cast<__half_raw>(f)), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat16 data type, relies on \p __NV_SATFINITE + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_bfloat16raw_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const __nv_bfloat16 f) { + __x = __nv_cvt_bfloat16raw_to_e8m0(static_cast<__nv_bfloat16_raw>(f), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float data type, relies on \p __NV_SATFINITE behavior + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_float_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const float f) { + __x = __nv_cvt_float_to_e8m0(f, __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double data type, relies on \p __NV_SATFINITE + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_double_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const double f) { + __x = __nv_cvt_double_to_e8m0(f, __NV_SATFINITE, cudaRoundPosInf); + } + + /* Converts from integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p short \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e8m0(const unsigned short int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const unsigned int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e8m0(const unsigned long long int val) { + __nv_bfloat16 rn = __ull2bfloat16_rn(val); + __nv_bfloat16_raw rn_raw = static_cast<__nv_bfloat16_raw>(rn); + unsigned long long int back_int = __bfloat162ull_rz(rn); + if (back_int < val) + { + rn_raw.x++; + } + __x = __nv_cvt_bfloat16raw_to_e8m0(rn_raw, __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const unsigned long int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p short \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const short int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const long long int val) { + __nv_bfloat16 rn = __ll2bfloat16_rn(val); + __nv_bfloat16_raw rn_raw = static_cast<__nv_bfloat16_raw>(rn); + long long int back_int = __bfloat162ll_rz(rn); + if (((val > 0) && (back_int < val)) || ((val < 0) && (back_int > val))) + { + rn_raw.x++; + } + __x = __nv_cvt_bfloat16raw_to_e8m0(rn_raw, __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const long int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening FP converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float() const { + return __internal_bf16raw_to_float(__nv_cvt_e8m0_to_bf16raw((*this).__x)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat16 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat16() const { + return static_cast<__nv_bfloat16>( + __nv_cvt_e8m0_to_bf16raw((*this).__x)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p double data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator double() const { + const float f = float(*this); + double d; +#if (defined __CUDA_ARCH__) + // protect from ftz in device code + asm("{ cvt.f64.f32 %0, %1;}\n" : "=d"(d) : "f"(f)); +#else + d = static_cast(f); +#endif + return d; + } + + /* rounding conversion to half */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half() const { + return __float2half_rn(float(*this)); + } + + /* Convert to integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned char() const { + unsigned char i; + const float f = float(*this); + const unsigned char max_val = 0xFFU; + const unsigned char bits = (*this).__x; + // saturation fixup + if (bits == 0xFFU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p short \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned short int() const { + return __bfloat162ushort_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned int() const { + return __bfloat162uint_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long int() const { + unsigned long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(unsigned long) == sizeof(unsigned long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__bfloat162ull_rz(__nv_bfloat16(*this))); + } + else + { + retval = static_cast(__bfloat162uint_rz(__nv_bfloat16(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000ULL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long long int() const { + return __bfloat162ull_rz(__nv_bfloat16(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p signed \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator signed char() const { + signed char i; + const float f = float(*this); + const signed char max_val = (signed char)0x7FU; + const unsigned char bits = (*this).__x; + // saturation fixup + if (bits == 0xFFU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to an implementation defined \p char data type. + * + * Detects signedness of the \p char type and proceeds accordingly, see + * further details in signed and unsigned char operators. + + * Clamps inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator char() const { + char value; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (((char)-1) < (char)0) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + value = static_cast(static_cast(*this)); + } + else + { + value = static_cast(static_cast(*this)); + } + return value; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p short \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator short int() const { + return __bfloat162short_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator int() const { + return __bfloat162int_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long int() const { + long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(long) == sizeof(long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__bfloat162ll_rz(__nv_bfloat16(*this))); + } + else + { + retval = static_cast(__bfloat162int_rz(__nv_bfloat16(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000LL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long long int() const { + return __bfloat162ll_rz(__nv_bfloat16(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p bool data type. + * All values in input range are non-zero, so result is always \p true. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator bool() const { + return true; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X2_E8M0_STRUCT C++ struct for handling vector type of two scale factors of e8m0 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X2_E8M0_STRUCT + * \brief __nv_fp8x2_e8m0 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of two scale factors of \p e8m0 kind each. + */ +struct __CUDA_ALIGN__(2) __nv_fp8x2_e8m0 { + public: + /** + * \ingroup CUDA_MATH_FP8X2_E8M0_STRUCT + * Storage variable contains the vector of two scale factor + * values. + */ + __nv_fp8x2_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x2_e8m0() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const __half2 f) { + __x = __nv_cvt_float2_to_e8m0x2(__half22float2(f), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat162 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const __nv_bfloat162 f) { + __x = __nv_cvt_bfloat162raw_to_e8m0x2(static_cast<__nv_bfloat162_raw>(f), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const float2 f) { + __x = __nv_cvt_float2_to_e8m0x2(f, + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const double2 f) { + __x = __nv_cvt_double2_to_e8m0x2(f, __NV_SATFINITE, cudaRoundPosInf); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat162 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat162() const { + return static_cast<__nv_bfloat162>(__nv_cvt_e8m0x2_to_bf162raw((*this).__x)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float2() const { + return __bfloat1622float2(static_cast<__nv_bfloat162>(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half2() const { + return __float22half2_rn(static_cast(*this)); + } + +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X4_E8M0_STRUCT C++ struct for handling vector type of four scale factors of e8m0 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X4_E8M0_STRUCT + * \brief __nv_fp8x4_e8m0 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of scale factors of \p e8m0 kind each. + */ +struct __CUDA_ALIGN__(4) __nv_fp8x4_e8m0 { + public: + /** + * \ingroup CUDA_MATH_FP8X4_E8M0_STRUCT + * Storage variable contains the vector of four scale factor + * values. + */ + __nv_fp8x4_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x4_e8m0() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __half2 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const __half2 flo, + const __half2 fhi) { + const __nv_fp8x2_storage_t rlo = static_cast<__nv_fp8x2_e8m0>(flo).__x; + const __nv_fp8x2_storage_t rhi = static_cast<__nv_fp8x2_e8m0>(fhi).__x; + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __nv_bfloat162 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const __nv_bfloat162 flo, + const __nv_bfloat162 fhi) { + const __nv_fp8x2_storage_t rlo = static_cast<__nv_fp8x2_e8m0>(flo).__x; + const __nv_fp8x2_storage_t rhi = static_cast<__nv_fp8x2_e8m0>(fhi).__x; + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const float4 f) { + const float2 flo = {f.x, f.y}; + const float2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_float2_to_e8m0x2(flo, __NV_SATFINITE, cudaRoundPosInf); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_float2_to_e8m0x2(fhi, __NV_SATFINITE, cudaRoundPosInf); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const double4 f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_e8m0x2(flo, __NV_SATFINITE, cudaRoundPosInf); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_e8m0x2(fhi, __NV_SATFINITE, cudaRoundPosInf); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float4 vector data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float4() const { + __nv_fp8x2_e8m0 lo; + lo.__x = static_cast<__nv_fp8x2_storage_t>(__x); + __nv_fp8x2_e8m0 hi; + hi.__x = static_cast<__nv_fp8x2_storage_t>(__x >> 16U); + + float2 rlo = static_cast(lo); + float2 rhi = static_cast(hi); + float4 res = {rlo.x, rlo.y, rhi.x, rhi.y}; + return res; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +#endif /* defined(__cplusplus) */ + +#endif /* end of include guard: __CUDA_FP8_HPP__ */ diff --git a/numba_cuda/numba/cuda/include/13/cuda_fp8.h b/numba_cuda/numba/cuda/include/13/cuda_fp8.h new file mode 100644 index 000000000..682190f44 --- /dev/null +++ b/numba_cuda/numba/cuda/include/13/cuda_fp8.h @@ -0,0 +1,496 @@ +/* + * Copyright 2022-2025 NVIDIA Corporation. All rights reserved. + * + * NOTICE TO LICENSEE: + * + * This source code and/or documentation ("Licensed Deliverables") are + * subject to NVIDIA intellectual property rights under U.S. and + * international Copyright laws. + * + * These Licensed Deliverables contained herein is PROPRIETARY and + * CONFIDENTIAL to NVIDIA and is being provided under the terms and + * conditions of a form of NVIDIA software license agreement by and + * between NVIDIA and Licensee ("License Agreement") or electronically + * accepted by Licensee. Notwithstanding any terms or conditions to + * the contrary in the License Agreement, reproduction or disclosure + * of the Licensed Deliverables to any third party without the express + * written consent of NVIDIA is prohibited. + * + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE + * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS + * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND. + * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED + * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY, + * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY + * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY + * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS, + * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS + * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE + * OF THESE LICENSED DELIVERABLES. + * + * U.S. Government End Users. These Licensed Deliverables are a + * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT + * 1995), consisting of "commercial computer software" and "commercial + * computer software documentation" as such terms are used in 48 + * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government + * only as a commercial end item. Consistent with 48 C.F.R.12.212 and + * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all + * U.S. Government End Users acquire the Licensed Deliverables with + * only those rights set forth herein. + * + * Any use of the Licensed Deliverables in individual and commercial + * software must include, in the user documentation and internal + * comments to the code, the above Disclaimer and U.S. Government End + * Users Notice. + */ + +#ifndef __CUDA_FP8_H__ +#define __CUDA_FP8_H__ + +/* Set up function decorations */ +#if defined(__CUDACC__) +#define __CUDA_FP8_DECL__ static __device__ __inline__ +#define __CUDA_HOSTDEVICE_FP8__ __host__ __device__ +#define __CUDA_HOSTDEVICE_FP8_DECL__ static __host__ __device__ __inline__ +#else /* !defined(__CUDACC__) */ +#if defined(__GNUC__) +#define __CUDA_HOSTDEVICE_FP8_DECL__ static __attribute__((unused)) +#else +#define __CUDA_HOSTDEVICE_FP8_DECL__ static +#endif /* defined(__GNUC__) */ +#define __CUDA_HOSTDEVICE_FP8__ +#endif /* defined(__CUDACC_) */ + +#if !defined(_MSC_VER) && __cplusplus >= 201103L +#define __CPP_VERSION_AT_LEAST_11_FP8 +#elif _MSC_FULL_VER >= 190024210 && _MSVC_LANG >= 201103L +#define __CPP_VERSION_AT_LEAST_11_FP8 +#endif + +#if ( (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000) && (defined __CUDA_ARCH_FAMILY_SPECIFIC__) ) + #if (( \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ >= 1000) && \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ < 1100) \ + ) || ( \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ >= 1100) && \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ < 1200) \ + ) || ( \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ >= 1200) && \ + (__CUDA_ARCH_FAMILY_SPECIFIC__ < 1300) \ + )) + #define __CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__ (1) + #else + // instructions unsupported, fallback to emulation + #define __CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__ (0) + #endif +#else /* ( (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000) && (defined __CUDA_ARCH_FAMILY_SPECIFIC__) ) */ + // instructions unsupported, fallback to emulation + #define __CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__ (0) +#endif /* ( (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000) && (defined __CUDA_ARCH_FAMILY_SPECIFIC__) ) */ + +// implicitly provided by NVRTC +#if !defined(__CUDACC_RTC__) +/* bring in enum cudaRoundMode */ +#include "device_types.h" +#endif /* !defined(__CUDACC_RTC__) */ + +/* bring in __half_raw data type */ +#include "cuda_fp16.h" +/* bring in __nv_bfloat16_raw data type */ +#include "cuda_bf16.h" + +// implicitly provided by NVRTC +#if !defined(__CUDACC_RTC__) +/* bring in float2, double4, etc vector types */ +#include "vector_types.h" +#endif /* !defined(__CUDACC_RTC__) */ + +/** + * \defgroup CUDA_MATH_INTRINSIC_FP8 FP8 Intrinsics + * This section describes fp8 intrinsic functions. + * To use these functions, include the header file \p cuda_fp8.h in your + * program. + * The following macros are available to help users selectively enable/disable + * various definitions present in the header file: + * - \p __CUDA_NO_FP8_CONVERSIONS__ - If defined, this macro will prevent any + * use of the C++ type conversions (converting constructors and conversion + * operators) defined in the header. + * - \p __CUDA_NO_FP8_CONVERSION_OPERATORS__ - If defined, this macro will + * prevent any use of the C++ conversion operators from \p fp8 to other types. + */ + +/** + * \defgroup CUDA_MATH_FP8_MISC FP8 Conversion and Data Movement + * \ingroup CUDA_MATH_INTRINSIC_FP8 + * To use these functions, include the header file \p cuda_fp8.h in your + * program. + */ + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief 8-bit \p unsigned \p integer + * type abstraction used for \p fp8 floating-point + * numbers storage. + */ +typedef unsigned char __nv_fp8_storage_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief 16-bit \p unsigned \p integer + * type abstraction used for storage of pairs of + * \p fp8 floating-point numbers. + */ +typedef unsigned short int __nv_fp8x2_storage_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief 32-bit \p unsigned \p integer + * type abstraction used for storage of tetrads of + * \p fp8 floating-point numbers. + */ +typedef unsigned int __nv_fp8x4_storage_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Enumerates the modes applicable when + * performing a narrowing conversion to \p fp8 destination types. + */ +typedef enum __nv_saturation_t { + /** + * Means no saturation to finite is performed when conversion + * results in rounding values outside the range of destination + * type. + * NOTE: for fp8 type of e4m3 kind, the results that are larger + * than the maximum representable finite number of the target + * format become NaN. + */ + __NV_NOSAT, + /** + * Means input larger than the maximum representable + * finite number MAXNORM of the target format round to the + * MAXNORM of the same sign as input. + */ + __NV_SATFINITE, +} __nv_saturation_t; + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Enumerates the possible + * interpretations of the 8-bit values when referring to them as + * \p fp8 types. + */ +typedef enum __nv_fp8_interpretation_t { + __NV_E4M3, /**< Stands for \p fp8 numbers of \p e4m3 kind. */ + __NV_E5M2, /**< Stands for \p fp8 numbers of \p e5m2 kind. */ +} __nv_fp8_interpretation_t; + +/* Forward-declaration of C-style APIs */ + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p double precision \p x to \p fp8 type of the + * requested kind using round-to-nearest-even rounding and requested saturation + * mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_double_to_fp8(const double x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p double precision numbers packed + * in \p double2 \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_double2_to_fp8x2(const double2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p single precision \p x to \p fp8 type of the + * requested kind using round-to-nearest-even rounding and requested saturation + * mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_float_to_fp8(const float x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p single precision numbers packed + * in \p float2 \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_float2_to_fp8x2(const float2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p half precision \p x to \p fp8 type of the requested + * kind using round-to-nearest-even rounding and requested saturation mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_halfraw_to_fp8(const __half_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p half precision numbers packed + * in \p __half2_raw \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_halfraw2_to_fp8x2( + const __half2_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p nv_bfloat16 precision \p x to \p fp8 type of the + * requested kind using round-to-nearest-even rounding and requested saturation + * mode. + * + * \details Converts input \p x to \p fp8 type of the kind specified by + * \p fp8_interpretation parameter, + * using round-to-nearest-even rounding and + * saturation mode specified by \p saturate parameter. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_fp8( + const __nv_bfloat16_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p nv_bfloat16 precision numbers packed + * in \p __nv_bfloat162_raw \p x into a vector of two values of \p fp8 type of + * the requested kind using round-to-nearest-even rounding and requested + * saturation mode. + * + * \details Converts input vector \p x to a vector of two \p fp8 values of the + * kind specified by \p fp8_interpretation parameter, using + * round-to-nearest-even rounding and saturation mode specified by \p saturate + * parameter. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_bfloat16raw2_to_fp8x2( + const __nv_bfloat162_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p fp8 \p x of the specified kind + * to \p half precision. + * + * \details Converts input \p x of \p fp8 type of the kind specified by + * \p fp8_interpretation parameter + * to \p half precision. + * + * \returns + * - The \p __half_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __half_raw +__nv_cvt_fp8_to_halfraw(const __nv_fp8_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation); +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input vector of two \p fp8 values of the specified kind + * to a vector of two \p half precision values packed in \p __half2_raw + * structure. + * + * \details Converts input vector \p x of \p fp8 type of the kind specified by + * \p fp8_interpretation parameter + * to a vector of two \p half precision values and returns as \p __half2_raw + * structure. + * + * \returns + * - The \p __half2_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __half2_raw +__nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation); + + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p bfloat16 input into a scaling factor of \p e8m0 kind. + * + * \details Input number's absolute value is rounded to the closest power of two in the + * direction specified via \p rounding parameter. Rounded results that are + * smaller than the smallest representable target format number 2^-127 are then + * clipped to 2^-127. Results that are larger than the largest representable + * target format number 2^127 are either clipped to 2^127 if \p saturate equals + * to \p __NV_SATFINITE, or convert to \p NaN otherwise. \p NaN inputs convert + * into \p NaN output, encoded as \p 0xFF in the target format. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_e8m0(const __nv_bfloat16_raw x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts a pair of \p bfloat16 values into a pair of scaling factors of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_bfloat162raw_to_e8m0x2(const __nv_bfloat162_raw x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p float value into a scaling factor of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_float_to_e8m0(const float x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts a pair of \p float values into a pair of scaling factors of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_float2_to_e8m0x2(const float2 x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input \p double value into a scaling factor of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_double_to_e8m0(const double x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts a pair of \p double values into a pair of scaling factors of \p e8m0 kind. + * + * \see __nv_cvt_bfloat16raw_to_e8m0() for details of conversion. + * + * \returns + * - The \p __nv_fp8x2_storage_t value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_double2_to_e8m0x2(const double2 x, const __nv_saturation_t saturate, const enum cudaRoundMode rounding); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input scaling factor value of \p e8m0 kind into \p bfloat16. + * + * \details Input scales are exact powers of two or a \p NaN value, + * also representable in the target format. + * + * \returns + * - The \p __nv_bfloat16_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat16_raw __nv_cvt_e8m0_to_bf16raw(const __nv_fp8_storage_t x); + +/** + * \ingroup CUDA_MATH_FP8_MISC + * \brief Converts input pair of scaling factors of \p e8m0 kind into a pair of \p bfloat16 values. + * + * \returns + * - The \p __nv_bfloat162_raw value holds the result of conversion. + */ +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat162_raw __nv_cvt_e8m0x2_to_bf162raw(const __nv_fp8x2_storage_t x); + +#if defined(__cplusplus) + +#define __CUDA_FP8_TYPES_EXIST__ + +/* Forward-declaration of structures defined in "cuda_fp8.hpp" */ +struct __nv_fp8_e5m2; +struct __nv_fp8x2_e5m2; +struct __nv_fp8x4_e5m2; + +struct __nv_fp8_e4m3; +struct __nv_fp8x2_e4m3; +struct __nv_fp8x4_e4m3; + +struct __nv_fp8_e8m0; +struct __nv_fp8x2_e8m0; +struct __nv_fp8x4_e8m0; + +#endif /* defined(__cplusplus) */ + +#include "cuda_fp8.hpp" + +#undef __CUDA_FP8_DECL__ +#undef __CUDA_HOSTDEVICE_FP8__ +#undef __CUDA_HOSTDEVICE_FP8_DECL__ + +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) +#undef __CPP_VERSION_AT_LEAST_11_FP8 +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#endif /* end of include guard: __CUDA_FP8_H__ */ diff --git a/numba_cuda/numba/cuda/include/13/cuda_fp8.hpp b/numba_cuda/numba/cuda/include/13/cuda_fp8.hpp new file mode 100644 index 000000000..6ee8c3d3d --- /dev/null +++ b/numba_cuda/numba/cuda/include/13/cuda_fp8.hpp @@ -0,0 +1,2804 @@ +/* + * Copyright 2022-2024 NVIDIA Corporation. All rights reserved. + * + * NOTICE TO LICENSEE: + * + * This source code and/or documentation ("Licensed Deliverables") are + * subject to NVIDIA intellectual property rights under U.S. and + * international Copyright laws. + * + * These Licensed Deliverables contained herein is PROPRIETARY and + * CONFIDENTIAL to NVIDIA and is being provided under the terms and + * conditions of a form of NVIDIA software license agreement by and + * between NVIDIA and Licensee ("License Agreement") or electronically + * accepted by Licensee. Notwithstanding any terms or conditions to + * the contrary in the License Agreement, reproduction or disclosure + * of the Licensed Deliverables to any third party without the express + * written consent of NVIDIA is prohibited. + * + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE + * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS + * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND. + * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED + * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY, + * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. + * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE + * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY + * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY + * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS, + * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS + * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE + * OF THESE LICENSED DELIVERABLES. + * + * U.S. Government End Users. These Licensed Deliverables are a + * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT + * 1995), consisting of "commercial computer software" and "commercial + * computer software documentation" as such terms are used in 48 + * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government + * only as a commercial end item. Consistent with 48 C.F.R.12.212 and + * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all + * U.S. Government End Users acquire the Licensed Deliverables with + * only those rights set forth herein. + * + * Any use of the Licensed Deliverables in individual and commercial + * software must include, in the user documentation and internal + * comments to the code, the above Disclaimer and U.S. Government End + * Users Notice. + */ + +#if !defined(__CUDA_FP8_HPP__) +#define __CUDA_FP8_HPP__ + +#if !defined(__CUDA_FP8_H__) +#error "Do not include this file directly. Instead, include cuda_fp8.h." +#endif + +/* C++ header for ::std::memcpy (used for type punning in host-side + * implementations). When compiling as a CUDA source file memcpy is provided + * implicitly. !defined(__CUDACC__) implies !defined(__CUDACC_RTC__). + */ +#if defined(__cplusplus) && !defined(__CUDACC__) +#include +#elif !defined(__cplusplus) && !defined(__CUDACC__) +#include +#endif /* defined(__cplusplus) && !defined(__CUDACC__) */ + +// implicitly provided by NVRTC +#if !defined(__CUDACC_RTC__) +#include +#endif /* !defined(__CUDACC_RTC__) */ + +#if !defined(IF_DEVICE_OR_CUDACC) +#if defined(__CUDACC__) + #define IF_DEVICE_OR_CUDACC(d, c, f) NV_IF_ELSE_TARGET(NV_IS_DEVICE, d, c) +#else + #define IF_DEVICE_OR_CUDACC(d, c, f) NV_IF_ELSE_TARGET(NV_IS_DEVICE, d, f) +#endif +#endif + +/* + * Bring in the standard assertions header to enforce the subset + * of rounding modes supported by the APIs defined here. + * NOTE: NVRTC defines its own assert + */ +#if !defined (__CUDACC_RTC__) +#include +#endif + +/* Set up structure-alignment attribute */ +#if !(defined __CUDA_ALIGN__) +#if defined(__CUDACC__) +#define __CUDA_ALIGN__(align) __align__(align) +#else +/* Define alignment macro based on compiler type (cannot assume C11 "_Alignas" + * is available) */ +#if __cplusplus >= 201103L +#define __CUDA_ALIGN__(n) \ + alignas(n) /* C++11 kindly gives us a keyword for this */ +#else /* !defined(__CPP_VERSION_AT_LEAST_11_FP8)*/ +#if defined(__GNUC__) +#define __CUDA_ALIGN__(n) __attribute__((aligned(n))) +#elif defined(_MSC_VER) +#define __CUDA_ALIGN__(n) __declspec(align(n)) +#else +#define __CUDA_ALIGN__(n) +#endif /* defined(__GNUC__) */ +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ +#endif /* defined(__CUDACC__) */ +#endif /* !(defined __CUDA_ALIGN__) */ + +#if !(defined __CPP_VERSION_AT_LEAST_11_FP8) +/* need c++11 for explicit operators */ +#define __CUDA_NO_FP8_CONVERSION_OPERATORS__ +#endif + +#if !(defined __DOXYGEN_ONLY__) + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_double_to_fp8(const double x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + unsigned char res; + unsigned long long int xbits; + +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&xbits, &x, sizeof(x)); +#else + (void)::std::memcpy(&xbits, &x, sizeof(x)); +#endif + unsigned char FP8_MAXNORM; + unsigned char FP8_MANTISSA_MASK; + unsigned short int FP8_EXP_BIAS; + unsigned long long int FP8_SIGNIFICAND_BITS; + const unsigned long long int DP_INF_BITS = 0x7FF0000000000000ULL; + unsigned long long int FP8_MINDENORM_O2; + unsigned long long int FP8_OVERFLOW_THRESHOLD; + unsigned long long int FP8_MINNORM; + + if (fp8_interpretation == __NV_E4M3) { + FP8_EXP_BIAS = 7U; + FP8_SIGNIFICAND_BITS = 4ULL; + FP8_MANTISSA_MASK = 0x7U; + FP8_MINDENORM_O2 = 0x3F50000000000000ULL; // mindenorm/2 = 2^-10 + FP8_OVERFLOW_THRESHOLD = + 0x407D000000000000ULL; // maxnorm + 1/2ulp = 0x1.Cp+8 + 0x1p+4 + FP8_MAXNORM = 0x7EU; + FP8_MINNORM = 0x3F90000000000000ULL; // minnorm = 2^-6 + } else { //__NV_E5M2 + FP8_EXP_BIAS = 15U; + FP8_SIGNIFICAND_BITS = 3ULL; + FP8_MANTISSA_MASK = 0x3U; + FP8_MINDENORM_O2 = 0x3EE0000000000000ULL; // mindenorm/2 = 2^-17 + FP8_OVERFLOW_THRESHOLD = + 0x40EE000000000000ULL - + 1ULL; // maxnorm + 1/2ulp = 0x1.Ep+15, and -1 to have common code + FP8_MAXNORM = 0x7BU; + FP8_MINNORM = 0x3F10000000000000ULL; // minnorm = 2^-14 + } + + // 1/2 LSB of the target format, positioned in double precision mantissa + // helpful in midpoints detection during round-to-nearest-even step + const unsigned long long int FP8_DP_HALF_ULP = + (unsigned long long int)1ULL << (53ULL - FP8_SIGNIFICAND_BITS - 1ULL); + // prepare sign bit in target format + unsigned char sign = (unsigned char)((xbits >> 63ULL) << 7U); + // prepare exponent field in target format + unsigned char exp = + (unsigned char)((((unsigned short int)(xbits >> 52ULL)) & 0x7FFU) - + 1023U + FP8_EXP_BIAS); + // round mantissa to target format width, rounding towards zero + unsigned char mantissa = + (unsigned char)(xbits >> (53ULL - FP8_SIGNIFICAND_BITS)) & + FP8_MANTISSA_MASK; + unsigned long long int absx = xbits & 0x7FFFFFFFFFFFFFFFULL; + + if (absx <= FP8_MINDENORM_O2) { + // zero or underflow + res = 0U; + } else if (absx > DP_INF_BITS) { + // NaN + if (fp8_interpretation == __NV_E4M3) { + res = 0x7FU; + } else { + // NaN --> QNaN + res = 0x7EU | mantissa; + } + } else if (absx > FP8_OVERFLOW_THRESHOLD) { + if (saturate == __NV_SATFINITE) { + res = FP8_MAXNORM; + } else { + // __NV_NOSAT + if (fp8_interpretation == __NV_E4M3) { + // no Inf in E4M3 + res = 0x7FU; // NaN + } else { + res = 0x7CU; // Inf in E5M2 + } + } + } else if (absx >= FP8_MINNORM) { + res = (unsigned char)((exp << (FP8_SIGNIFICAND_BITS - 1U)) | mantissa); + // rounded-off bits + unsigned long long int round = + xbits & ((FP8_DP_HALF_ULP << 1ULL) - 1ULL); + // round-to-nearest-even adjustment + if ((round > FP8_DP_HALF_ULP) || + ((round == FP8_DP_HALF_ULP) && (mantissa & 1U))) { + res = (unsigned char)(res + 1U); + } + } else // Denormal range + { + unsigned char shift = (unsigned char)(1U - exp); + // add implicit leading bit + mantissa |= (unsigned char)(1U << (FP8_SIGNIFICAND_BITS - 1U)); + // additional round-off due to denormalization + res = (unsigned char)(mantissa >> shift); + + // rounded-off bits, including implicit leading bit + unsigned long long int round = + (xbits | ((unsigned long long int)1ULL << (53ULL - 1ULL))) & + ((FP8_DP_HALF_ULP << (shift + 1ULL)) - 1ULL); + // round-to-nearest-even adjustment + if ((round > (FP8_DP_HALF_ULP << shift)) || + ((round == (FP8_DP_HALF_ULP << shift)) && (res & 1U))) { + res = (unsigned char)(res + 1U); + } + } + + res |= sign; + + return (__nv_fp8_storage_t)res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_double2_to_fp8x2(const double2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8x2_storage_t storage = (__nv_fp8x2_storage_t)__nv_cvt_double_to_fp8( + x.y, saturate, fp8_interpretation); + storage = (__nv_fp8x2_storage_t)(storage << 8U); + storage = (__nv_fp8x2_storage_t)(storage | + __nv_cvt_double_to_fp8( + x.x, saturate, fp8_interpretation)); + return storage; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_float_to_fp8(const float x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8_storage_t res = 0U; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + __nv_fp8x2_storage_t storage; + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x), "f"(0.0f)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x), "f"(0.0f)); + } + res = (__nv_fp8_storage_t)storage; + } else +#endif + { + unsigned int xbits; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&xbits, &x, sizeof(x)); +#else + (void)::std::memcpy(&xbits, &x, sizeof(x)); +#endif + + // isnan + if ((xbits & 0x7FFFFFFFU) > 0x7F800000U) { + // Canonical NaN + xbits = 0x7FFFFFFFU; + } + + float fx; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&fx, &xbits, sizeof(xbits)); +#else + (void)::std::memcpy(&fx, &xbits, sizeof(xbits)); +#endif + + const double dx = (double)fx; + res = __nv_cvt_double_to_fp8(dx, saturate, fp8_interpretation); + } + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_float2_to_fp8x2(const float2 x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8x2_storage_t storage; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x.x), "f"(x.y)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f32 %0, %2, %1;}\n" + : "=h"(storage) + : "f"(x.x), "f"(x.y)); + } + } else +#endif + { + storage = (__nv_fp8x2_storage_t)__nv_cvt_float_to_fp8( + x.y, saturate, fp8_interpretation); + storage = (__nv_fp8x2_storage_t)(storage << 8U); + storage = (__nv_fp8x2_storage_t)(storage | __nv_cvt_float_to_fp8( + x.x, saturate, + fp8_interpretation)); + } + return storage; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float +__internal_halfraw_to_float(const __half_raw x) { + float f; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 530) + asm("{cvt.f32.f16 %0, %1;}\n" : "=f"(f) : "h"(x.x)); +#else + const unsigned int ux = (unsigned int)x.x; + unsigned int sign = (ux >> 15U) & 1U; + unsigned int exponent = (ux >> 10U) & 0x1fU; + unsigned int mantissa = (ux & 0x3ffU) << 13U; + if (exponent == 0x1fU) { /* NaN or Inf */ + /* discard sign of a NaN */ + sign = ((mantissa != 0U) ? (sign >> 1U) : sign); + mantissa = ((mantissa != 0U) ? 0x7fffffU : 0U); + exponent = 0xffU; + } else if (exponent == 0U) { /* Denorm or Zero */ + if (mantissa != 0U) { + unsigned int msb; + exponent = 0x71U; + do { + msb = (mantissa & 0x400000U); + mantissa <<= 1U; /* normalize */ + --exponent; + } while (msb == 0U); + mantissa &= 0x7fffffU; /* 1.mantissa is implicit */ + } + } else { + exponent += 0x70U; + } + const unsigned int u = ((sign << 31U) | (exponent << 23U) | mantissa); +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&f, &u, sizeof(u)); +#else + (void)::std::memcpy(&f, &u, sizeof(u)); +#endif +#endif /* (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 530) */ + return f; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float2 +__internal_halfraw2_to_float2(const __half2_raw x) { + __half_raw raw; + float2 res; + raw.x = x.x; + res.x = __internal_halfraw_to_float(raw); + raw.x = x.y; + res.y = __internal_halfraw_to_float(raw); + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t +__nv_cvt_halfraw_to_fp8(const __half_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8_storage_t res = 0U; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + unsigned int half2_storage = (unsigned int)(x.x); + __nv_fp8x2_storage_t tmp; + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } + res = (__nv_fp8_storage_t)tmp; + } else +#endif + { + float fx = __internal_halfraw_to_float(x); + res = __nv_cvt_float_to_fp8(fx, saturate, fp8_interpretation); + } + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t __nv_cvt_halfraw2_to_fp8x2( + const __half2_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_fp8x2_storage_t tmp; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + if (saturate == __NV_SATFINITE) { + unsigned int half2_storage; + (void)memcpy(&half2_storage, &x, sizeof(x)); + + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.satfinite.e5m2x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } else { + asm("{cvt.rn.satfinite.e4m3x2.f16x2 %0, %1;}\n" + : "=h"(tmp) + : "r"(half2_storage)); + } + } else +#endif + { + __half_raw raw; + raw.x = x.x; + __nv_fp8_storage_t lo = + __nv_cvt_halfraw_to_fp8(raw, saturate, fp8_interpretation); + raw.x = x.y; + __nv_fp8_storage_t hi = + __nv_cvt_halfraw_to_fp8(raw, saturate, fp8_interpretation); + tmp = hi; + tmp = (__nv_fp8x2_storage_t)(tmp << 8U); + tmp = (__nv_fp8x2_storage_t)(tmp | lo); + } + return tmp; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float +__internal_bf16raw_to_float(const __nv_bfloat16_raw x) { + const unsigned int ux = ((unsigned int)x.x) << 16U; + float fx; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&fx, &ux, sizeof(ux)); +#else + (void)::std::memcpy(&fx, &ux, sizeof(ux)); +#endif + return fx; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8_storage_t __nv_cvt_bfloat16raw_to_fp8( + const __nv_bfloat16_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + const float fx = __internal_bf16raw_to_float(x); + const __nv_fp8_storage_t res = + __nv_cvt_float_to_fp8(fx, saturate, fp8_interpretation); + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_fp8x2_storage_t +__nv_cvt_bfloat16raw2_to_fp8x2( + const __nv_bfloat162_raw x, const __nv_saturation_t saturate, + const __nv_fp8_interpretation_t fp8_interpretation) { + __nv_bfloat16_raw raw; + raw.x = x.y; + __nv_fp8x2_storage_t storage = + (__nv_fp8x2_storage_t)__nv_cvt_bfloat16raw_to_fp8(raw, saturate, + fp8_interpretation); + storage = (__nv_fp8x2_storage_t)(storage << 8U); + raw.x = x.x; + storage = (__nv_fp8x2_storage_t)(storage | + __nv_cvt_bfloat16raw_to_fp8( + raw, saturate, fp8_interpretation)); + return storage; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __half2_raw +__nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation); + +__CUDA_HOSTDEVICE_FP8_DECL__ __half_raw +__nv_cvt_fp8_to_halfraw(const __nv_fp8_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation) { + __half_raw res; + res.x = 0U; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + res.x = + __nv_cvt_fp8x2_to_halfraw2((__nv_fp8x2_storage_t)x, fp8_interpretation) + .x; +#else + unsigned short int ur = (unsigned short int)x; + ur = (unsigned short int)(ur << 8U); + + if (fp8_interpretation == __NV_E5M2) { + if ((ur & 0x7FFFU) > 0x7C00U) { + /* If NaN, return canonical NaN */ + ur = 0x7FFFU; + } + } else { // __NV_E4M3 + unsigned short int sign = ur & 0x8000U; + unsigned short int exponent = + (unsigned short int)(((ur & 0x7800U) >> 1U) + 0x2000U); + unsigned short int mantissa = (ur & 0x0700U) >> 1U; + unsigned char absx = 0x7FU & (unsigned char)x; + + if (absx == 0x7FU) // NaN + { + ur = 0x7FFFU; // fp16 canonical NaN, discard sign + } else if (exponent == 0x2000U) { + // zero or denormal + if (mantissa != 0U) { + // normalize + mantissa = (unsigned short int)(mantissa << 1U); + while ((mantissa & 0x0400U) == 0U) { + mantissa = (unsigned short int)(mantissa << 1U); + exponent = (unsigned short int)(exponent - 0x0400U); + } + // discard implicit leading bit + mantissa &= 0x03FFU; + } else { // Zero + exponent = 0U; + } + + ur = (sign | exponent) | mantissa; + } else { + ur = (sign | exponent) | mantissa; + } + } + res.x = ur; +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __half2_raw +__nv_cvt_fp8x2_to_halfraw2(const __nv_fp8x2_storage_t x, + const __nv_fp8_interpretation_t fp8_interpretation) { + __half2_raw res; +#if (defined __CUDA_ARCH__) && (__CUDA_ARCH__ >= 890) + unsigned int half2_storage; + if (fp8_interpretation == __NV_E5M2) { + asm("{cvt.rn.f16x2.e5m2x2 %0, %1;}\n" : "=r"(half2_storage) : "h"(x)); + } else { + asm("{cvt.rn.f16x2.e4m3x2 %0, %1;}\n" : "=r"(half2_storage) : "h"(x)); + } + (void)memcpy(&res, &half2_storage, sizeof(half2_storage)); +#else + res.x = + __nv_cvt_fp8_to_halfraw((__nv_fp8_storage_t)x, fp8_interpretation).x; + res.y = __nv_cvt_fp8_to_halfraw((__nv_fp8_storage_t)(x >> 8U), + fp8_interpretation) + .x; +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8_storage_t __nv_cvt_bfloat16raw_to_e8m0(const __nv_bfloat16_raw x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short ures = 0U; + unsigned in = (unsigned)(x.x); + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } + return (__nv_fp8_storage_t)ures; +#else + // extract exponent bits, provides non-saturated result in RZ + __nv_fp8_storage_t res = (unsigned char)(x.x >> 7U); + + if (rounding == cudaRoundPosInf) { + // round-up if mantissa non-zero and |x| > 2^-127 and finite + if ((x.x & 0x007FU) && ((x.x & 0x7FFFU) > 0x0040U) && ((x.x & 0x7FFFU) < 0x7F80U)) res++; + } + + // Handle saturation of non-NaN large inputs to finite + if (saturate == __NV_SATFINITE) { + // non-NaN, Overflow --> Max + if (((x.x & 0x7FFFU) <= 0x7F80U) && (res == 0xFFU)) + { + res--; + } + } + return res; +#endif +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8x2_storage_t __nv_cvt_bfloat162raw_to_e8m0x2(const __nv_bfloat162_raw x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + assert((rounding == cudaRoundZero) || (rounding == cudaRoundPosInf)); +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + __nv_fp8x2_storage_t ures = 0U; + unsigned in = (unsigned)(x.x) | ((unsigned)(x.y) << (unsigned)16U); + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.bf16x2 %0, %1;}\n" + : "=h"(ures) + : "r"(in)); + } + return ures; +#else + __nv_bfloat16_raw lo, hi; + lo.x = x.x; + hi.x = x.y; + __nv_fp8x2_storage_t ures = __nv_cvt_bfloat16raw_to_e8m0(hi, saturate, rounding); + ures <<= (unsigned short)8U; + ures |= __nv_cvt_bfloat16raw_to_e8m0(lo, saturate, rounding); + return ures; +#endif +} + +__CUDA_HOSTDEVICE_FP8_DECL__ unsigned int __internal_fp8_float_as_uint(const float f) +{ + unsigned int u; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&u, &f, sizeof(f)); +#else + (void)::std::memcpy(&u, &f, sizeof(f)); +#endif + return u; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ float __internal_fp8_uint_as_float(const unsigned int u) +{ + float f; +#if defined(__CUDACC__) || (!defined __cplusplus) + (void)memcpy(&f, &u, sizeof(u)); +#else + (void)::std::memcpy(&f, &u, sizeof(u)); +#endif + return f; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat16_raw +__internal_float_to_bf16raw_rz(const float x) { + __nv_bfloat16_raw r; +NV_IF_ELSE_TARGET(NV_PROVIDES_SM_80, + asm("{ cvt.rz.bf16.f32 %0, %1;}\n" : "=h"(r.x) : "f"(x)); +, + unsigned int ux = __internal_fp8_float_as_uint(x); + if ((ux & 0x7FFFFFFFU) > 0x7f800000U) + { + // NaN + r.x = (unsigned short int)0x7FFFU; + } + else + { + r.x = (unsigned short int)(ux >> 16U); + } +) + return r; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ __nv_bfloat16_raw +__internal_float_to_bf16raw_ru(const float x) { + __nv_bfloat16_raw r; +NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90, + asm("{ cvt.rp.bf16.f32 %0, %1;}\n" : "=h"(r.x) : "f"(x)); +, + unsigned int ux = __internal_fp8_float_as_uint(x); + if ((ux & 0x7FFFFFFFU) > 0x7f800000U) + { + // NaN + r.x = (unsigned short int)0x7FFFU; + } + else if ((ux < 0x7f800000U) && ((ux & 0x0000FFFFU) != 0)) + { + // 0 <= x < +inf, round-up + r.x = (unsigned short int)((ux >> 16U) + 1U); + } + else { + // truncate others + r.x = (unsigned short int)(ux >> 16U); + } +) + return r; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8_storage_t __nv_cvt_float_to_e8m0(const float x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + assert((rounding == cudaRoundZero) || (rounding == cudaRoundPosInf)); + __nv_fp8_storage_t res = 0U; +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short ures = 0U; + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(ures) + : "f"(x), "f"(0.0f)); + } + res = (__nv_fp8_storage_t)ures; +#else + if (rounding == cudaRoundZero) + { + res = __nv_cvt_bfloat16raw_to_e8m0(__internal_float_to_bf16raw_rz(x), saturate, rounding); + } + else + { //cudaRoundPosInf + float absx = __internal_fp8_uint_as_float((__internal_fp8_float_as_uint(x) << (unsigned)1U) >> (unsigned)1U); + res = __nv_cvt_bfloat16raw_to_e8m0(__internal_float_to_bf16raw_ru(absx), saturate, rounding); + } +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8x2_storage_t __nv_cvt_float2_to_e8m0x2(const float2 x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + assert((rounding == cudaRoundZero) || (rounding == cudaRoundPosInf)); + __nv_fp8x2_storage_t res = 0U; +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + if ((rounding == cudaRoundZero) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rz.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } else if ((rounding == cudaRoundZero) && (saturate == __NV_NOSAT)) { + asm("{cvt.rz.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_SATFINITE)) { + asm("{cvt.rp.satfinite.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } else if ((rounding == cudaRoundPosInf) && (saturate == __NV_NOSAT)) { + asm("{cvt.rp.ue8m0x2.f32 %0, %2, %1;}\n" + : "=h"(res) + : "f"(x.x), "f"(x.y)); + } +#else + res = __nv_cvt_float_to_e8m0(x.y, saturate, rounding); + res <<= (unsigned short)8U; + res |= __nv_cvt_float_to_e8m0(x.x, saturate, rounding); +#endif + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +float __internal_double_to_float_with_sticky(double x) +{ +#if (defined __CUDA_ARCH__) + // protect from ftz in device code + float f; + double d; + asm("{ cvt.rn.f32.f64 %0, %1;}\n" : "=f"(f) : "d"(x)); + asm("{ cvt.f64.f32 %0, %1;}\n" : "=d"(d) : "f"(f)); +#else + const float f = (float)x; + const double d = (double)f; +#endif + unsigned int u = __internal_fp8_float_as_uint(f); + int x_is_not_nan = ((u << (unsigned)1U) <= (unsigned)0xFF000000U) ? 1 : 0; + + if ((x > 0.0) && (d > x)) { + u--; + } + if ((x < 0.0) && (d < x)) { + u--; + } + if ((d != x) && (x_is_not_nan == 1)) { + u |= 1U; + } + return __internal_fp8_uint_as_float(u); +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8_storage_t __nv_cvt_double_to_e8m0(const double x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + float fx_with_sticky = __internal_double_to_float_with_sticky(x); + __nv_fp8_storage_t res = __nv_cvt_float_to_e8m0(fx_with_sticky, saturate, rounding); + + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_fp8x2_storage_t __nv_cvt_double2_to_e8m0x2(const double2 x, + const __nv_saturation_t saturate, const enum cudaRoundMode rounding) +{ + float2 f; + f.x = __internal_double_to_float_with_sticky(x.x); + f.y = __internal_double_to_float_with_sticky(x.y); + return __nv_cvt_float2_to_e8m0x2(f, saturate, rounding); +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +unsigned short __internal_e8m0_to_bf16(const __nv_fp8_storage_t x) +{ + unsigned short res; + // shift bias exponent bits into place + res = ((unsigned short)x) << 7U; + + if (x == 0xFFU) { + res = 0x7FFFU; // NaN --> Canonical QNaN + } else if (x == 0U) { + res = 0x0040U; // 2^-127 + } + + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_bfloat16_raw __nv_cvt_e8m0_to_bf16raw(const __nv_fp8_storage_t x) +{ + __nv_bfloat16_raw res; + +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short in = (unsigned short)x; + unsigned hr = 0U; + asm("{cvt.rn.bf16x2.ue8m0x2 %0, %1;}\n" + : "=r"(hr) + : "h"(in)); + + res.x = (unsigned short)hr; +#else + res.x = __internal_e8m0_to_bf16(x); +#endif + + return res; +} + +__CUDA_HOSTDEVICE_FP8_DECL__ +__nv_bfloat162_raw __nv_cvt_e8m0x2_to_bf162raw(const __nv_fp8x2_storage_t x) +{ + __nv_bfloat162_raw res; + +#if (__CUDA_FP8_INTERNAL_CAN_RELY_ON_PTX_FOR_SHORTTYPESCVT__) + unsigned short in = (unsigned short)x; + unsigned hr = 0U; + asm("{cvt.rn.bf16x2.ue8m0x2 %0, %1;}\n" + : "=r"(hr) + : "h"(in)); + + res.x = (unsigned short)hr; + res.y = (unsigned short)(hr >> (unsigned)16U); +#else + + res.x = __internal_e8m0_to_bf16((__nv_fp8_storage_t)x); + res.y = __internal_e8m0_to_bf16((__nv_fp8_storage_t)(x >> (unsigned short)8U)); + +#endif + + return res; +} + +#endif /* !(defined __DOXYGEN_ONLY__) */ + +/* All other definitions in this file are only visible to C++ compilers */ +#if defined(__cplusplus) + +/** + * \defgroup CUDA_MATH_FP8_E5M2_STRUCT C++ struct for handling fp8 data type of e5m2 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8_E5M2_STRUCT + * \brief __nv_fp8_e5m2 datatype + * + * \details This structure implements the datatype for handling + * \p fp8 floating-point numbers of \p e5m2 kind: + * with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(1) __nv_fp8_e5m2 { + public: + /** + * \ingroup CUDA_MATH_FP8_E5M2_STRUCT + * Storage variable contains the \p fp8 floating-point data. + */ + __nv_fp8_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8_e5m2() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider FP types */ + /* Note we do avoid constructor init-list because of special host/device + * compilation rules */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const __half f) { + __x = __nv_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat16 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const __nv_bfloat16 f) { + __x = __nv_cvt_bfloat16raw_to_fp8(static_cast<__nv_bfloat16_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const float f) { + __x = __nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const double f) { + __x = __nv_cvt_double_to_fp8(f, __NV_SATFINITE, __NV_E5M2); + } + + /* Converts from integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p short \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e5m2(const unsigned short int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const unsigned int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const unsigned long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p long \p int data type, relies on + * \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e5m2(const unsigned long long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p short \p int data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const short int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p int data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p int data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e5m2(const long long int val) { + __x = static_cast<__nv_fp8_e5m2>(static_cast(val)).__x; + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening FP converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half() const { + return static_cast<__half>(__nv_cvt_fp8_to_halfraw(__x, __NV_E5M2)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float() const { + return __internal_halfraw_to_float( + __nv_cvt_fp8_to_halfraw(__x, __NV_E5M2)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat16 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat16() const { + return __float2bfloat16_rz(float(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p double data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator double() const { + return static_cast(float(*this)); + } + + /* Convert to integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p char data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned char() const { + unsigned char i; + const float f = float(*this); + const unsigned char max_val = 0xFFU; + const unsigned char min_val = 0U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) > 0x7CU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p short \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned short int() const { + return __half2ushort_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned int() const { + return __half2uint_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long int() const { + unsigned long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(unsigned long) == sizeof(unsigned long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ull_rz(__half(*this))); + } + else + { + retval = static_cast(__half2uint_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p long \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000ULL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long long int() const { + return __half2ull_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p signed \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator signed char() const { + signed char i; + const float f = float(*this); + const signed char max_val = (signed char)0x7FU; + const signed char min_val = (signed char)0x80U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) > 0x7CU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to an implementation defined \p char data type. + * + * Detects signedness of the \p char type and proceeds accordingly, see + * further details in signed and unsigned char operators. + + * Clamps inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator char() const { + char value; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (((char)-1) < (char)0) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + value = static_cast(static_cast(*this)); + } + else + { + value = static_cast(static_cast(*this)); + } + return value; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p short \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator short int() const { + return __half2short_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator int() const { + return __half2int_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long int() const { + long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(long) == sizeof(long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ll_rz(__half(*this))); + } + else + { + retval = static_cast(__half2int_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000LL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long long int() const { + return __half2ll_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p bool data type. + * +0 and -0 inputs convert to \p false. + * Non-zero inputs convert to \p true. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator bool() const { + return (__x & 0x7FU) != 0U; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X2_E5M2_STRUCT C++ struct for handling vector type of two fp8 values of e5m2 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X2_E5M2_STRUCT + * \brief __nv_fp8x2_e5m2 datatype + * + * \details This structure implements the datatype for handling two + * \p fp8 floating-point numbers of \p e5m2 kind each: + * with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(2) __nv_fp8x2_e5m2 { + public: + /** + * \ingroup CUDA_MATH_FP8X2_E5M2_STRUCT + * Storage variable contains the vector of two \p fp8 floating-point data + * values. + */ + __nv_fp8x2_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x2_e5m2() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const __half2 f) { + __x = __nv_cvt_halfraw2_to_fp8x2(static_cast<__half2_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat162 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const __nv_bfloat162 f) { + __x = __nv_cvt_bfloat16raw2_to_fp8x2(static_cast<__nv_bfloat162_raw>(f), + __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const float2 f) { + __x = __nv_cvt_float2_to_fp8x2(f, __NV_SATFINITE, __NV_E5M2); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e5m2(const double2 f) { + __x = __nv_cvt_double2_to_fp8x2(f, __NV_SATFINITE, __NV_E5M2); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half2() const { + return static_cast<__half2>(__nv_cvt_fp8x2_to_halfraw2(__x, __NV_E5M2)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float2() const { + return __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(__x, __NV_E5M2)); + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +__CUDA_HOSTDEVICE_FP8_DECL__ unsigned int +__internal_pack_u16x2_to_u32(const unsigned short int src_lo, + const unsigned short int src_hi) { + unsigned int dst; +#if (defined __CUDACC__) && (defined __CUDA_ARCH__) + asm("{ mov.b32 %0, {%1,%2};}\n" : "=r"(dst) : "h"(src_lo), "h"(src_hi)); +#else + dst = (static_cast(src_hi) << 16U) | + static_cast(src_lo); +#endif + return dst; +} + +/** + * \defgroup CUDA_MATH_FP8X4_E5M2_STRUCT C++ struct for handling vector type of four fp8 values of e5m2 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X4_E5M2_STRUCT + * \brief __nv_fp8x4_e5m2 datatype + * + * \details This structure implements the datatype for handling four + * \p fp8 floating-point numbers of \p e5m2 kind each: + * with 1 sign, 5 exponent, 1 implicit and 2 explicit mantissa bits. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(4) __nv_fp8x4_e5m2 { + public: + /** + * \ingroup CUDA_MATH_FP8X4_E5M2_STRUCT + * Storage variable contains the vector of four \p fp8 floating-point data + * values. + */ + __nv_fp8x4_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x4_e5m2() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __half2 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const __half2 flo, + const __half2 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(flo), __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(fhi), __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __nv_bfloat162 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const __nv_bfloat162 flo, + const __nv_bfloat162 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(flo), __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(fhi), __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const float4 f) { + const float2 flo = {f.x, f.y}; + const float2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_float2_to_fp8x2(flo, __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_float2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + __NV_SILENCE_DEPRECATION_BEGIN + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const double4 f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + __NV_SILENCE_DEPRECATION_END + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4_16a vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const double4_16a f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4_32a vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e5m2(const double4_32a f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E5M2); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E5M2); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float4 vector data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float4() const { + const __nv_fp8x2_storage_t slo = static_cast<__nv_fp8x2_storage_t>(__x); + const __nv_fp8x2_storage_t shi = + static_cast<__nv_fp8x2_storage_t>(__x >> 16U); + float2 rlo = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(slo, __NV_E5M2)); + float2 rhi = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(shi, __NV_E5M2)); + float4 res = {rlo.x, rlo.y, rhi.x, rhi.y}; + return res; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8_E4M3_STRUCT C++ struct for handling fp8 data type of e4m3 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8_E4M3_STRUCT + * \brief __nv_fp8_e4m3 datatype + * + * \details This structure implements the datatype for storing + * \p fp8 floating-point numbers of \p e4m3 kind: + * with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. + * The encoding doesn't support Infinity. + * NaNs are limited to 0x7F and 0xFF values. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(1) __nv_fp8_e4m3 { + public: + /** + * \ingroup CUDA_MATH_FP8_E4M3_STRUCT + * Storage variable contains the \p fp8 floating-point data. + */ + __nv_fp8_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8_e4m3() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider FP types */ + /* Note we do avoid constructor init-list because of special host/device + * compilation rules */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const __half f) { + __x = __nv_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat16 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const __nv_bfloat16 f) { + __x = __nv_cvt_bfloat16raw_to_fp8(static_cast<__nv_bfloat16_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const float f) { + __x = __nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const double f) { + __x = __nv_cvt_double_to_fp8(f, __NV_SATFINITE, __NV_E4M3); + } + + /* Converts from integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p short \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e4m3(const unsigned short int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const unsigned int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const unsigned long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p long \p int data type, relies on + * \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e4m3(const unsigned long long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p short \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const short int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p int data type, relies on \p __NV_SATFINITE behavior + * for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p long \p int data type, relies on \p + * __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e4m3(const long long int val) { + __x = static_cast<__nv_fp8_e4m3>(static_cast(val)).__x; + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening FP converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half() const { + return static_cast<__half>(__nv_cvt_fp8_to_halfraw(__x, __NV_E4M3)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float() const { + return __internal_halfraw_to_float( + __nv_cvt_fp8_to_halfraw(__x, __NV_E4M3)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat16 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat16() const { + return __float2bfloat16_rz(float(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p double data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator double() const { + return static_cast(float(*this)); + } + + /* Convert to integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p char data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned char() const { + unsigned char i; + const float f = float(*this); + const unsigned char max_val = 0xFFU; + const unsigned char min_val = 0U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) == 0x7FU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p short \p int data type. + * Clamps negative inputs to zero. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned short int() const { + return __half2ushort_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p int data type. + * Clamps negative inputs to zero. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned int() const { + return __half2uint_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p int data type. + * Clamps negative and too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long int() const { + unsigned long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(unsigned long) == sizeof(unsigned long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ull_rz(__half(*this))); + } + else + { + retval = static_cast(__half2uint_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p long \p int data type. + * Clamps negative inputs to zero. + * \p NaN inputs convert to \p 0x8000000000000000ULL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long long int() const { + return __half2ull_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p signed \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator signed char() const { + signed char i; + const float f = float(*this); + const signed char max_val = (signed char)0x7FU; + const signed char min_val = (signed char)0x80U; + const unsigned char bits = (*this).__x; + // saturation fixup + if ((bits & 0x7FU) == 0x7FU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else if (f < static_cast(min_val)) { + // saturate minimum + i = min_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to an implementation defined \p char data type. + * + * Detects signedness of the \p char type and proceeds accordingly, see + * further details in signed and unsigned char operators. + + * Clamps inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator char() const { + char value; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (((char)-1) < (char)0) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + value = static_cast(static_cast(*this)); + } + else + { + value = static_cast(static_cast(*this)); + } + return value; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p short \p int data type. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator short int() const { + return __half2short_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p int data type. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator int() const { + return __half2int_rz(__half(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long int() const { + long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(long) == sizeof(long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__half2ll_rz(__half(*this))); + } + else + { + retval = static_cast(__half2int_rz(__half(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p long \p int data type. + * \p NaN inputs convert to \p 0x8000000000000000LL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long long int() const { + return __half2ll_rz(__half(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p bool data type. + * +0 and -0 inputs convert to \p false. + * Non-zero inputs convert to \p true. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator bool() const { + return (__x & 0x7FU) != 0U; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X2_E4M3_STRUCT C++ struct for handling vector type of two fp8 values of e4m3 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X2_E4M3_STRUCT + * \brief __nv_fp8x2_e4m3 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of two \p fp8 values of \p e4m3 kind each: + * with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. + * The encoding doesn't support Infinity. + * NaNs are limited to 0x7F and 0xFF values. + */ +struct __CUDA_ALIGN__(2) __nv_fp8x2_e4m3 { + public: + /** + * \ingroup CUDA_MATH_FP8X2_E4M3_STRUCT + * Storage variable contains the vector of two \p fp8 floating-point data + * values. + */ + __nv_fp8x2_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x2_e4m3() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const __half2 f) { + __x = __nv_cvt_halfraw2_to_fp8x2(static_cast<__half2_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat162 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const __nv_bfloat162 f) { + __x = __nv_cvt_bfloat16raw2_to_fp8x2(static_cast<__nv_bfloat162_raw>(f), + __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const float2 f) { + __x = __nv_cvt_float2_to_fp8x2(f, __NV_SATFINITE, __NV_E4M3); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e4m3(const double2 f) { + __x = __nv_cvt_double2_to_fp8x2(f, __NV_SATFINITE, __NV_E4M3); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half2() const { + return static_cast<__half2>(__nv_cvt_fp8x2_to_halfraw2(__x, __NV_E4M3)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float2() const { + return __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(__x, __NV_E4M3)); + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X4_E4M3_STRUCT C++ struct for handling vector type of four fp8 values of e4m3 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X4_E4M3_STRUCT + * \brief __nv_fp8x4_e4m3 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of four \p fp8 values of \p e4m3 kind each: + * with 1 sign, 4 exponent, 1 implicit and 3 explicit mantissa bits. + * The encoding doesn't support Infinity. + * NaNs are limited to 0x7F and 0xFF values. + */ +struct __CUDA_ALIGN__(4) __nv_fp8x4_e4m3 { + public: + /** + * \ingroup CUDA_MATH_FP8X4_E4M3_STRUCT + * Storage variable contains the vector of four \p fp8 floating-point data + * values. + */ + __nv_fp8x4_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x4_e4m3() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __half2 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const __half2 flo, + const __half2 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(flo), __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = __nv_cvt_halfraw2_to_fp8x2( + static_cast<__half2_raw>(fhi), __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __nv_bfloat162 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const __nv_bfloat162 flo, + const __nv_bfloat162 fhi) { + const __nv_fp8x2_storage_t rlo = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(flo), __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = __nv_cvt_bfloat16raw2_to_fp8x2( + static_cast<__nv_bfloat162_raw>(fhi), __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const float4 f) { + const float2 flo = {f.x, f.y}; + const float2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_float2_to_fp8x2(flo, __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_float2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + __NV_SILENCE_DEPRECATION_BEGIN + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const double4 f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + __NV_SILENCE_DEPRECATION_END + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4_16a vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const double4_16a f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4_32a vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e4m3(const double4_32a f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_fp8x2(flo, __NV_SATFINITE, __NV_E4M3); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_fp8x2(fhi, __NV_SATFINITE, __NV_E4M3); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float4 vector data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float4() const { + const __nv_fp8x2_storage_t slo = static_cast<__nv_fp8x2_storage_t>(__x); + const __nv_fp8x2_storage_t shi = + static_cast<__nv_fp8x2_storage_t>(__x >> 16U); + float2 rlo = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(slo, __NV_E4M3)); + float2 rhi = __internal_halfraw2_to_float2( + __nv_cvt_fp8x2_to_halfraw2(shi, __NV_E4M3)); + float4 res = {rlo.x, rlo.y, rhi.x, rhi.y}; + return res; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \ingroup CUDA_MATH_FP8_E8M0_STRUCT + * \brief __nv_fp8_e8m0 datatype + * + * \details This structure implements the datatype for handling + * 8-bit scale factors of \p e8m0 kind: interpreted as powers of two + * with biased exponent. Bias equals to 127, so numbers 0 through 254 + * represent 2^-127 through 2^127. Number \p 0xFF = 255 is reserved + * for NaN. + * + * The structure implements converting constructors and operators. + */ +struct __CUDA_ALIGN__(1) __nv_fp8_e8m0 { + public: + /** + * \ingroup CUDA_MATH_FP8_E8M0_STRUCT + * Storage variable contains the 8-bit scale data. + */ + __nv_fp8_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8_e8m0() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider FP types */ + /* Note we do avoid constructor init-list because of special host/device + * compilation rules */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half data type, relies on \p __NV_SATFINITE + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_float_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const __half f) { + __x = __nv_cvt_float_to_e8m0(__internal_halfraw_to_float(static_cast<__half_raw>(f)), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat16 data type, relies on \p __NV_SATFINITE + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_bfloat16raw_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const __nv_bfloat16 f) { + __x = __nv_cvt_bfloat16raw_to_e8m0(static_cast<__nv_bfloat16_raw>(f), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float data type, relies on \p __NV_SATFINITE behavior + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_float_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const float f) { + __x = __nv_cvt_float_to_e8m0(f, __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double data type, relies on \p __NV_SATFINITE + * behavior for large input values and \p cudaRoundPosInf for + * rounding. + * \see __nv_cvt_double_to_e8m0 for further details + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const double f) { + __x = __nv_cvt_double_to_e8m0(f, __NV_SATFINITE, cudaRoundPosInf); + } + + /* Converts from integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p short \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e8m0(const unsigned short int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const unsigned int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ + __nv_fp8_e8m0(const unsigned long long int val) { + __nv_bfloat16 rn = __ull2bfloat16_rn(val); + __nv_bfloat16_raw rn_raw = static_cast<__nv_bfloat16_raw>(rn); + unsigned long long int back_int = __bfloat162ull_rz(rn); + if (back_int < val) + { + rn_raw.x++; + } + __x = __nv_cvt_bfloat16raw_to_e8m0(rn_raw, __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p unsigned \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const unsigned long int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p short \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const short int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const long long int val) { + __nv_bfloat16 rn = __ll2bfloat16_rn(val); + __nv_bfloat16_raw rn_raw = static_cast<__nv_bfloat16_raw>(rn); + long long int back_int = __bfloat162ll_rz(rn); + if (((val > 0) && (back_int < val)) || ((val < 0) && (back_int > val))) + { + rn_raw.x++; + } + __x = __nv_cvt_bfloat16raw_to_e8m0(rn_raw, __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p long \p int data type, relies on + * \p cudaRoundPosInf rounding. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8_e8m0(const long int val) { + __x = static_cast<__nv_fp8_e8m0>(static_cast(val)).__x; + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening FP converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float() const { + return __internal_bf16raw_to_float(__nv_cvt_e8m0_to_bf16raw((*this).__x)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat16 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat16() const { + return static_cast<__nv_bfloat16>( + __nv_cvt_e8m0_to_bf16raw((*this).__x)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p double data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator double() const { + const float f = float(*this); + double d; +#if (defined __CUDA_ARCH__) + // protect from ftz in device code + asm("{ cvt.f64.f32 %0, %1;}\n" : "=d"(d) : "f"(f)); +#else + d = static_cast(f); +#endif + return d; + } + + /* rounding conversion to half */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half() const { + return __float2half_rn(float(*this)); + } + + /* Convert to integral */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned char() const { + unsigned char i; + const float f = float(*this); + const unsigned char max_val = 0xFFU; + const unsigned char bits = (*this).__x; + // saturation fixup + if (bits == 0xFFU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p short \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned short int() const { + return __bfloat162ushort_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned int() const { + return __bfloat162uint_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long int() const { + unsigned long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(unsigned long) == sizeof(unsigned long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__bfloat162ull_rz(__nv_bfloat16(*this))); + } + else + { + retval = static_cast(__bfloat162uint_rz(__nv_bfloat16(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p unsigned \p long \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000ULL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator unsigned long long int() const { + return __bfloat162ull_rz(__nv_bfloat16(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p signed \p char data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator signed char() const { + signed char i; + const float f = float(*this); + const signed char max_val = (signed char)0x7FU; + const unsigned char bits = (*this).__x; + // saturation fixup + if (bits == 0xFFU) { + // NaN + i = 0; + } else if (f > static_cast(max_val)) { + // saturate maximum + i = max_val; + } else { + // normal value + i = static_cast(f); + } + return i; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to an implementation defined \p char data type. + * + * Detects signedness of the \p char type and proceeds accordingly, see + * further details in signed and unsigned char operators. + + * Clamps inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator char() const { + char value; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (((char)-1) < (char)0) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + value = static_cast(static_cast(*this)); + } + else + { + value = static_cast(static_cast(*this)); + } + return value; + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p short \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator short int() const { + return __bfloat162short_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator int() const { + return __bfloat162int_rz(__nv_bfloat16(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p zero if output type is 32-bit. + * \p NaN inputs convert to \p 0x8000000000000000ULL if output type is 64-bit. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long int() const { + long retval; + /* Suppress VS warning: warning C4127: conditional expression is constant */ +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (push) +#pragma warning (disable: 4127) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + if (sizeof(long) == sizeof(long long)) +#if defined(_MSC_VER) && !defined(__CUDA_ARCH__) +#pragma warning (pop) +#endif /* _MSC_VER && !defined(__CUDA_ARCH__) */ + { + retval = static_cast(__bfloat162ll_rz(__nv_bfloat16(*this))); + } + else + { + retval = static_cast(__bfloat162int_rz(__nv_bfloat16(*this))); + } + return retval; + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p long \p long \p int data type. + * Clamps too large inputs to the output range. + * \p NaN inputs convert to \p 0x8000000000000000LL. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator long long int() const { + return __bfloat162ll_rz(__nv_bfloat16(*this)); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p bool data type. + * All values in input range are non-zero, so result is always \p true. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator bool() const { + return true; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X2_E8M0_STRUCT C++ struct for handling vector type of two scale factors of e8m0 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X2_E8M0_STRUCT + * \brief __nv_fp8x2_e8m0 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of two scale factors of \p e8m0 kind each. + */ +struct __CUDA_ALIGN__(2) __nv_fp8x2_e8m0 { + public: + /** + * \ingroup CUDA_MATH_FP8X2_E8M0_STRUCT + * Storage variable contains the vector of two scale factor + * values. + */ + __nv_fp8x2_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x2_e8m0() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __half2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const __half2 f) { + __x = __nv_cvt_float2_to_e8m0x2(__half22float2(f), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p __nv_bfloat162 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const __nv_bfloat162 f) { + __x = __nv_cvt_bfloat162raw_to_e8m0x2(static_cast<__nv_bfloat162_raw>(f), + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const float2 f) { + __x = __nv_cvt_float2_to_e8m0x2(f, + __NV_SATFINITE, cudaRoundPosInf); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double2 data type, relies on \p __NV_SATFINITE + * behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x2_e8m0(const double2 f) { + __x = __nv_cvt_double2_to_e8m0x2(f, __NV_SATFINITE, cudaRoundPosInf); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __nv_bfloat162 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __nv_bfloat162() const { + return static_cast<__nv_bfloat162>(__nv_cvt_e8m0x2_to_bf162raw((*this).__x)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float2() const { + return __bfloat1622float2(static_cast<__nv_bfloat162>(*this)); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p __half2 data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator __half2() const { + return __float22half2_rn(static_cast(*this)); + } + +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +/** + * \defgroup CUDA_MATH_FP8X4_E8M0_STRUCT C++ struct for handling vector type of four scale factors of e8m0 kind. + * \ingroup CUDA_MATH_INTRINSIC_FP8 + */ + +/** + * \ingroup CUDA_MATH_FP8X4_E8M0_STRUCT + * \brief __nv_fp8x4_e8m0 datatype + * + * \details This structure implements the datatype for storage + * and operations on the vector of scale factors of \p e8m0 kind each. + */ +struct __CUDA_ALIGN__(4) __nv_fp8x4_e8m0 { + public: + /** + * \ingroup CUDA_MATH_FP8X4_E8M0_STRUCT + * Storage variable contains the vector of four scale factor + * values. + */ + __nv_fp8x4_storage_t __x; + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor by default. + */ +#if defined(__CPP_VERSION_AT_LEAST_11_FP8) + __nv_fp8x4_e8m0() = default; +#else + __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0() {} +#endif /* defined(__CPP_VERSION_AT_LEAST_11_FP8) */ + +#if !defined(__CUDA_NO_FP8_CONVERSIONS__) + + /* Construct from wider types */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __half2 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const __half2 flo, + const __half2 fhi) { + const __nv_fp8x2_storage_t rlo = static_cast<__nv_fp8x2_e8m0>(flo).__x; + const __nv_fp8x2_storage_t rhi = static_cast<__nv_fp8x2_e8m0>(fhi).__x; + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from a pair of \p __nv_bfloat162 data type values, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const __nv_bfloat162 flo, + const __nv_bfloat162 fhi) { + const __nv_fp8x2_storage_t rlo = static_cast<__nv_fp8x2_e8m0>(flo).__x; + const __nv_fp8x2_storage_t rhi = static_cast<__nv_fp8x2_e8m0>(fhi).__x; + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p float4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const float4 f) { + const float2 flo = {f.x, f.y}; + const float2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_float2_to_e8m0x2(flo, __NV_SATFINITE, cudaRoundPosInf); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_float2_to_e8m0x2(fhi, __NV_SATFINITE, cudaRoundPosInf); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4 vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + __NV_SILENCE_DEPRECATION_BEGIN + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const double4 f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_e8m0x2(flo, __NV_SATFINITE, cudaRoundPosInf); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_e8m0x2(fhi, __NV_SATFINITE, cudaRoundPosInf); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + __NV_SILENCE_DEPRECATION_END + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4_16a vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const double4_16a f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_e8m0x2(flo, __NV_SATFINITE, cudaRoundPosInf); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_e8m0x2(fhi, __NV_SATFINITE, cudaRoundPosInf); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Constructor from \p double4_32a vector data type, + * relies on \p __NV_SATFINITE behavior for out-of-range values. + */ + explicit __CUDA_HOSTDEVICE_FP8__ __nv_fp8x4_e8m0(const double4_32a f) { + const double2 flo = {f.x, f.y}; + const double2 fhi = {f.z, f.w}; + const __nv_fp8x2_storage_t rlo = + __nv_cvt_double2_to_e8m0x2(flo, __NV_SATFINITE, cudaRoundPosInf); + const __nv_fp8x2_storage_t rhi = + __nv_cvt_double2_to_e8m0x2(fhi, __NV_SATFINITE, cudaRoundPosInf); + __x = __internal_pack_u16x2_to_u32(rlo, rhi); + } + +#if !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) + /* Widening converts */ + + /** + * \ingroup CUDA_MATH_FP8_MISC + * Conversion operator to \p float4 vector data type. + */ + explicit __CUDA_HOSTDEVICE_FP8__ operator float4() const { + __nv_fp8x2_e8m0 lo; + lo.__x = static_cast<__nv_fp8x2_storage_t>(__x); + __nv_fp8x2_e8m0 hi; + hi.__x = static_cast<__nv_fp8x2_storage_t>(__x >> 16U); + + float2 rlo = static_cast(lo); + float2 rhi = static_cast(hi); + float4 res = {rlo.x, rlo.y, rhi.x, rhi.y}; + return res; + } +#endif /* !defined(__CUDA_NO_FP8_CONVERSION_OPERATORS__) */ +#endif /* !defined(__CUDA_NO_FP8_CONVERSIONS__) */ +}; + +#endif /* defined(__cplusplus) */ + +#endif /* end of include guard: __CUDA_FP8_HPP__ */ diff --git a/numba_cuda/numba/cuda/simulator/api.py b/numba_cuda/numba/cuda/simulator/api.py index de070f310..81dff259a 100644 --- a/numba_cuda/numba/cuda/simulator/api.py +++ b/numba_cuda/numba/cuda/simulator/api.py @@ -37,6 +37,11 @@ def is_bfloat16_supported(): return False +def is_fp8_supported(): + """FP8 is not hardware-accelerated in the simulator.""" + return False + + class stream: """ The stream API is supported in the simulator - however, all execution @@ -81,6 +86,7 @@ def detect(): print("Found 1 CUDA devices") print("id %d %20s %40s" % (0, "SIMULATOR", "[SUPPORTED]")) print("%40s: 5.0" % "compute capability") + print("%40s: No" % "FP8 Hardware Acceleration") def list_devices(): diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index cbe39946a..387eea09b 100644 --- a/numba_cuda/numba/cuda/target.py +++ b/numba_cuda/numba/cuda/target.py @@ -43,6 +43,7 @@ def load_additional_registries(self): cudamath, fp16, bf16, + fp8, libdevicedecl, vector_types, ) @@ -58,6 +59,7 @@ def load_additional_registries(self): self.install_registry(vector_types.typing_registry) self.install_registry(fp16.typing_registry) self.install_registry(bf16.typing_registry) + self.install_registry(fp8.typing_registry) def resolve_value_type(self, val): # treat other dispatcher object as another device function @@ -186,6 +188,7 @@ def load_additional_registries(self): mathimpl as cuda_mathimpl, vector_types, bf16, + fp8, ) # fix for #8940 @@ -203,6 +206,7 @@ def load_additional_registries(self): self.install_registry(vector_types.impl_registry) self.install_registry(fp16.target_registry) self.install_registry(bf16.target_registry) + self.install_registry(fp8.target_registry) self.install_registry(slicing.registry) self.install_registry(iterators.registry) self.install_registry(listobj.registry) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py b/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py new file mode 100644 index 000000000..c3e9719fb --- /dev/null +++ b/numba_cuda/numba/cuda/tests/cudapy/test_fp8_bindings.py @@ -0,0 +1,564 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause + +""" +Comprehensive test suite for CUDA FP8 types and conversion functions. + +Tests cover: +- fp8_e5m2, fp8_e4m3, fp8_e8m0 types +- Constructors from various numeric types +- Conversion operators to various types +- Conversion intrinsics +""" + +import unittest +from numba import cuda +from numba.cuda.testing import unittest, CUDATestCase +import numpy as np + +from numba.cuda import ( + int8, + int16, + int32, + int64, + uint8, + uint16, + uint32, + uint64, + float32, + float64, +) +from numba.cuda.types import float16, bfloat16 +from numba.cuda import config +from numba.cuda.api import is_fp8_supported + +if not config.ENABLE_CUDASIM: + from numba.cuda._internal.cuda_fp8 import ( + fp8_e5m2, + fp8_e4m3, + fp8_e8m0, + cvt_float_to_fp8, + cvt_double_to_fp8, + cvt_bfloat16raw_to_fp8, + cvt_bfloat16raw_to_e8m0, + cvt_float_to_e8m0, + cvt_double_to_e8m0, + cvt_e8m0_to_bf16raw, + saturation_t, + fp8_interpretation_t, + ) + from cuda.bindings.runtime import cudaRoundMode + +FE8_TYPES = [fp8_e5m2, fp8_e4m3, fp8_e8m0] + + +@unittest.skipUnless(is_fp8_supported(), "FP8 is not supported") +@unittest.skipIf( + config.ENABLE_CUDASIM, "FP8 is not supported on CUDA simulator" +) +class FP8ConstructorTests(CUDATestCase): + """Basic constructor for FP8 types.""" + + def test_fp8_e5m2_constructors(self): + """Test fp8_e5m2 construction from all numeric types (floats and ints).""" + + @cuda.jit + def kernel(result): + # Test fp8_e5m2 from all float types + # Use 1.0 which is exactly representable in all FP8 formats + result[0] = float32(fp8_e5m2(float16(1.0))) + result[1] = float32(fp8_e5m2(float32(1.0))) + result[2] = float32(fp8_e5m2(float64(1.0))) + result[3] = float32(fp8_e5m2(bfloat16(1.0))) + + # Test fp8_e5m2 from all integer types + # Use 2 which is exactly representable (power of 2) + result[4] = float32(fp8_e5m2(int8(2))) + result[5] = float32(fp8_e5m2(int16(2))) + result[6] = float32(fp8_e5m2(int32(2))) + result[7] = float32(fp8_e5m2(int64(2))) + result[8] = float32(fp8_e5m2(uint8(2))) + result[9] = float32(fp8_e5m2(uint16(2))) + result[10] = float32(fp8_e5m2(uint32(2))) + result[11] = float32(fp8_e5m2(uint64(2))) + + result = np.zeros(12, dtype=np.float32) + kernel[1, 1](result) + + # Check float conversions (indices 0-3) - should be exactly 1.0 + for i in range(4): + self.assertEqual(result[i], 1.0) + + # Check integer conversions (indices 4-11) - should be exactly 2.0 + for i in range(4, 12): + self.assertEqual(result[i], 2.0) + + def test_fp8_e4m3_constructors(self): + """Test fp8_e4m3 construction from all numeric types (floats and ints).""" + + @cuda.jit + def kernel(result): + # Test fp8_e4m3 from all float types + # Use 1.0 which is exactly representable in all FP8 formats + result[0] = float32(fp8_e4m3(float16(1.0))) + result[1] = float32(fp8_e4m3(float32(1.0))) + result[2] = float32(fp8_e4m3(float64(1.0))) + result[3] = float32(fp8_e4m3(bfloat16(1.0))) + + # Test fp8_e4m3 from all integer types + # Use 2 which is exactly representable (power of 2) + result[4] = float32(fp8_e4m3(int8(2))) + result[5] = float32(fp8_e4m3(int16(2))) + result[6] = float32(fp8_e4m3(int32(2))) + result[7] = float32(fp8_e4m3(int64(2))) + result[8] = float32(fp8_e4m3(uint8(2))) + result[9] = float32(fp8_e4m3(uint16(2))) + result[10] = float32(fp8_e4m3(uint32(2))) + result[11] = float32(fp8_e4m3(uint64(2))) + + result = np.zeros(12, dtype=np.float32) + kernel[1, 1](result) + + # Check float conversions (indices 0-3) - should be exactly 1.0 + for i in range(4): + self.assertEqual(result[i], 1.0) + + # Check integer conversions (indices 4-11) - should be exactly 2.0 + for i in range(4, 12): + self.assertEqual(result[i], 2.0) + + def test_fp8_e8m0_constructors(self): + """Test fp8_e8m0 construction from all numeric types (floats and ints).""" + + @cuda.jit + def kernel(result): + # Test fp8_e8m0 from all float types (use power of 2 for e8m0) + result[0] = float32(fp8_e8m0(float16(2.0))) + result[1] = float32(fp8_e8m0(float32(2.0))) + result[2] = float32(fp8_e8m0(float64(2.0))) + result[3] = float32(fp8_e8m0(bfloat16(2.0))) + + # Test fp8_e8m0 from all integer types (use power of 2) + result[4] = float32(fp8_e8m0(int8(4))) + result[5] = float32(fp8_e8m0(int16(4))) + result[6] = float32(fp8_e8m0(int32(4))) + result[7] = float32(fp8_e8m0(int64(4))) + result[8] = float32(fp8_e8m0(uint8(4))) + result[9] = float32(fp8_e8m0(uint16(4))) + result[10] = float32(fp8_e8m0(uint32(4))) + result[11] = float32(fp8_e8m0(uint64(4))) + + result = np.zeros(12, dtype=np.float32) + kernel[1, 1](result) + + # Check float conversions - should be exactly 2.0 (e8m0 stores powers of 2) + for i in range(4): + self.assertEqual(result[i], 2.0) + + # Check integer conversions - should be exactly 4.0 + for i in range(4, 12): + self.assertEqual(result[i], 4.0) + + def test_fp8_nan_constructors(self): + """Test fp8 construction from NaN.""" + + @cuda.jit + def kernel(result): + nan = float32(float("nan")) + result[0] = float32(fp8_e5m2(nan)) + result[1] = float32(fp8_e4m3(nan)) + result[2] = float32(fp8_e8m0(nan)) + + result = np.zeros(3, dtype=np.float32) + kernel[1, 1](result) + + self.assertTrue(np.isnan(result[0])) + self.assertTrue(np.isnan(result[1])) + self.assertTrue(np.isnan(result[2])) + + +@unittest.skipUnless(is_fp8_supported(), "FP8 is not supported") +@unittest.skipIf( + config.ENABLE_CUDASIM, "FP8 is not supported on CUDA simulator" +) +class FP8ConversionTests(CUDATestCase): + """Test FP8 conversion operators to various types.""" + + def test_fp8_to_float_types(self): + """Test FP8 conversion to float types (__half, float, double, bfloat16).""" + + # Test data: (fp8_type, test_value, expected, tolerance_places) + test_cases = [ + (fp8_e5m2, 1.5, 1.5, 2), + (fp8_e4m3, 2.5, 2.5, 2), + (fp8_e8m0, 8.0, 8.0, None), # Power of 2, exact + ] + + for fp8_type, test_val, expected, places in test_cases: + with self.subTest(fp8_type=fp8_type.__name__, value=test_val): + + @cuda.jit + def kernel(result): + fp8_val = fp8_type(float32(test_val)) + + result[0] = float32(float16(fp8_val)) # to __half + result[1] = float32(fp8_val) # to float + result[2] = float32(float64(fp8_val)) # to double + result[3] = float32(bfloat16(fp8_val)) # to bfloat16 + + result = np.zeros(4, dtype=np.float32) + kernel[1, 1](result) + + if places is None: + self.assertTrue(np.all(result == expected)) + else: + self.assertTrue( + np.allclose( + result, expected, rtol=0, atol=10 ** (-places) + ) + ) + + def test_fp8_to_unsigned_integers(self): + """Test FP8 conversion to unsigned integer types.""" + + # Test data: (fp8_type, test_value, expected) + test_cases = [ + (fp8_e5m2, 10.0, 10), + (fp8_e4m3, 12.0, 12), + (fp8_e8m0, 16.0, 16), + ] + + for fp8_type, test_val, expected in test_cases: + with self.subTest(fp8_type=fp8_type.__name__, value=test_val): + + @cuda.jit + def kernel(result): + fp8_val = fp8_type(float32(test_val)) + + result[0] = uint64(uint8(fp8_val)) + result[1] = uint64(uint16(fp8_val)) + result[2] = uint64(uint32(fp8_val)) + result[3] = uint64(fp8_val) + + result = np.zeros(4, dtype=np.uint64) + kernel[1, 1](result) + + self.assertTrue(np.all(result == expected)) + + def test_fp8_to_signed_integers(self): + """Test FP8 conversion to signed integer types.""" + + # Test data: (fp8_type, positive_value, negative_value) + test_cases = [ + (fp8_e5m2, 16.0, -8.0), + (fp8_e4m3, 20.0, -6.0), + ( + fp8_e8m0, + 32.0, + 16.0, + ), # Because fp8_e8m0 is an exponent only type, it does not represent negative values. + ] + + for fp8_type, pos_val, neg_val in test_cases: + with self.subTest(fp8_type=fp8_type.__name__): + + @cuda.jit + def kernel(result): + fp8_pos = fp8_type(float32(pos_val)) + fp8_neg = fp8_type(float32(neg_val)) + + result[0] = int64(int8(fp8_pos)) + result[1] = int64(int16(fp8_pos)) + result[2] = int64(int32(fp8_pos)) + result[3] = int64(fp8_pos) + result[4] = int64(int8(fp8_neg)) + result[5] = int64(int16(fp8_neg)) + result[6] = int64(int32(fp8_neg)) + result[7] = int64(fp8_neg) + + result = np.zeros(8, dtype=np.int64) + kernel[1, 1](result) + + # Check positive conversions + np.testing.assert_array_equal( + result[:4], np.array([int(pos_val)] * 4) + ) + + # Check negative conversions + np.testing.assert_array_equal( + result[4:], np.array([int(neg_val)] * 4) + ) + + def test_fp8_conversion_edge_cases_zero(self): + """Test conversion of zero values for all FP8 types.""" + + for fp8_type in FE8_TYPES: + with self.subTest(fp8_type=fp8_type.__name__): + + @cuda.jit + def kernel(result): + zero = fp8_type(float32(0.0)) + result[0] = float32(zero) + result[1] = int64(zero) + + result = np.zeros(2, dtype=np.int64) + kernel[1, 1](result) + + np.testing.assert_array_equal(result, np.array([0, 0])) + + def test_fp8_conversion_negative_values(self): + """Test conversion of negative values for all FP8 types. + Because fp8_e8m0 is an exponent only type, it does not represent negative values. + """ + + # Test data: (fp8_type, test_value) + test_cases = [(fp8_e5m2, -3.0), (fp8_e4m3, -4.0)] + + for fp8_type, neg_val in test_cases: + with self.subTest(fp8_type=fp8_type.__name__, value=neg_val): + + @cuda.jit + def kernel(result_float, result_int): + fp8_neg = fp8_type(float32(neg_val)) + + result_float[0] = float32(fp8_neg) + result_int[0] = int64(int32(fp8_neg)) + + result_float = np.zeros(1, dtype=np.float32) + result_int = np.zeros(1, dtype=np.int64) + kernel[1, 1](result_float, result_int) + + self.assertAlmostEqual(result_float[0], neg_val, places=1) + self.assertAlmostEqual(result_int[0], int(neg_val), delta=1) + + def test_fp8_conversion_roundtrip(self): + """Test roundtrip conversions: fp8 -> float -> fp8 -> float.""" + + for fp8_type in FE8_TYPES: + with self.subTest(fp8_type=fp8_type.__name__): + + @cuda.jit + def kernel(result): + float_val = float32(4.0) + fp8_val2 = fp8_type(float_val) + result[0] = float32(fp8_val2) + + result = np.zeros(1, dtype=np.float32) + kernel[1, 1](result) + + np.testing.assert_array_equal(result, np.array([4.0])) + + def test_fp8_nan_conversions(self): + """Test FP8 NaN conversion to other types.""" + + @cuda.jit + def kernel(result): + nan = float32(float("nan")) + # Create FP8 NaNs + nan_e5m2 = fp8_e5m2(nan) + nan_e4m3 = fp8_e4m3(nan) + nan_e8m0 = fp8_e8m0(nan) + + # Convert back to float types + result[0] = float32(float16(nan_e5m2)) + result[1] = float32(nan_e5m2) + result[2] = float32(float64(nan_e5m2)) + + result[3] = float32(float16(nan_e4m3)) + result[4] = float32(nan_e4m3) + result[5] = float32(float64(nan_e4m3)) + + result[6] = float32(float16(nan_e8m0)) + result[7] = float32(nan_e8m0) + result[8] = float32(float64(nan_e8m0)) + + result = np.zeros(9, dtype=np.float32) + kernel[1, 1](result) + + np.testing.assert_array_equal(result, np.array([float("nan")] * 9)) + + +@unittest.skipUnless(is_fp8_supported(), "FP8 is not supported") +@unittest.skipIf( + config.ENABLE_CUDASIM, "FP8 is not supported on CUDA simulator" +) +class FP8Storage_CVT_Intrinsics_Tests(CUDATestCase): + """Test raw conversion intrinsics operating on storage types.""" + + def test_cvt_float_to_fp8(self): + @cuda.jit + def kernel(result, x): + # Use an out-of-range value so NOSAT and SATFINITE differ. + # For overflow: + # - E5M2: NOSAT -> Inf (0x7C), SATFINITE -> MAXNORM (0x7B) + # - E4M3: NOSAT -> NaN (0x7F), SATFINITE -> MAXNORM (0x7E) + result[0] = cvt_float_to_fp8( + x[0], saturation_t.NOSAT, fp8_interpretation_t.E5M2 + ) + result[1] = cvt_float_to_fp8( + x[0], saturation_t.SATFINITE, fp8_interpretation_t.E5M2 + ) + result[2] = cvt_float_to_fp8( + x[0], saturation_t.NOSAT, fp8_interpretation_t.E4M3 + ) + result[3] = cvt_float_to_fp8( + x[0], saturation_t.SATFINITE, fp8_interpretation_t.E4M3 + ) + + result = np.zeros(4, dtype=np.uint8) + x = np.array([1e20], dtype=np.float32) + kernel[1, 1](result, x) + + self.assertEqual(result[0], 0x7C) # E5M2 overflow -> Inf (NOSAT) + self.assertEqual( + result[1], 0x7B + ) # E5M2 overflow -> MAXNORM (SATFINITE) + self.assertEqual(result[2], 0x7F) # E4M3 overflow -> NaN (NOSAT) + self.assertEqual( + result[3], 0x7E + ) # E4M3 overflow -> MAXNORM (SATFINITE) + + def test_cvt_double_to_fp8(self): + @cuda.jit + def kernel(result, x): + result[0] = cvt_double_to_fp8( + x[0], saturation_t.NOSAT, fp8_interpretation_t.E5M2 + ) + result[1] = cvt_double_to_fp8( + x[0], saturation_t.SATFINITE, fp8_interpretation_t.E5M2 + ) + result[2] = cvt_double_to_fp8( + x[0], saturation_t.NOSAT, fp8_interpretation_t.E4M3 + ) + result[3] = cvt_double_to_fp8( + x[0], saturation_t.SATFINITE, fp8_interpretation_t.E4M3 + ) + + result = np.zeros(4, dtype=np.uint8) + x = np.array([1e300], dtype=np.float64) + kernel[1, 1](result, x) + self.assertEqual(result[0], 0x7C) # E5M2 overflow -> Inf (NOSAT) + self.assertEqual( + result[1], 0x7B + ) # E5M2 overflow -> MAXNORM (SATFINITE) + self.assertEqual(result[2], 0x7F) # E4M3 overflow -> NaN (NOSAT) + self.assertEqual( + result[3], 0x7E + ) # E4M3 overflow -> MAXNORM (SATFINITE) + + def test_cvt_e8m0_to_bf16raw(self): + @cuda.jit + def kernel(result, x): + raw = cvt_e8m0_to_bf16raw(x[0]) + result[0] = raw.x + + result = np.zeros(1, dtype=np.uint16) + # 1.0 in E8M0 is 127 (bias 127) + x = np.array([127], dtype=np.uint8) + kernel[1, 1](result, x) + + # 1.0 in BF16 is 0x3F80 + self.assertEqual(result[0], 0x3F80) + + def test_cvt_bfloat16raw_roundtrip(self): + @cuda.jit + def kernel(result, x): + # x is uint8 (e8m0) + # Convert e8m0 to bfloat16_raw + raw = cvt_e8m0_to_bf16raw(x[0]) + + # Convert bf16_raw to fp8 using both NOSAT and SATFINITE. + result[0] = cvt_bfloat16raw_to_fp8( + raw, saturation_t.NOSAT, fp8_interpretation_t.E5M2 + ) + result[1] = cvt_bfloat16raw_to_fp8( + raw, saturation_t.SATFINITE, fp8_interpretation_t.E5M2 + ) + result[2] = cvt_bfloat16raw_to_fp8( + raw, saturation_t.NOSAT, fp8_interpretation_t.E4M3 + ) + result[3] = cvt_bfloat16raw_to_fp8( + raw, saturation_t.SATFINITE, fp8_interpretation_t.E4M3 + ) + + result = np.zeros(4, dtype=np.uint8) + # 2^127 in E8M0 (very out-of-range for fp8) + x = np.array([254], dtype=np.uint8) + kernel[1, 1](result, x) + + self.assertEqual(result[0], 0x7C) # E5M2 overflow -> Inf (NOSAT) + self.assertEqual( + result[1], 0x7B + ) # E5M2 overflow -> MAXNORM (SATFINITE) + self.assertEqual(result[2], 0x7F) # E4M3 overflow -> NaN (NOSAT) + self.assertEqual( + result[3], 0x7E + ) # E4M3 overflow -> MAXNORM (SATFINITE) + + 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) + + self.assertEqual(result[0], 0xFF) # NOSAT overflow -> NaN + self.assertEqual( + result[1], 0xFE + ) # SATFINITE overflow -> max finite (2^127) + + 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) + + self.assertEqual(result[0], 0xFF) # NOSAT overflow -> NaN + self.assertEqual( + result[1], 0xFE + ) # SATFINITE overflow -> max finite (2^127) + + 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) + + self.assertEqual(result[0], 126) + self.assertEqual(result[1], 126) + + +if __name__ == "__main__": + unittest.main() diff --git a/pyproject.toml b/pyproject.toml index 00ffce12f..fc355e95c 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -110,6 +110,7 @@ exclude = [ "*.h", "cuda_bf16.py", "cuda_fp16.py", + "cuda_fp8.py", ] [tool.ruff.lint.per-file-ignores]