diff --git a/configs/cuda_fp16.yml b/configs/cuda_fp16.yml index ba87c6a2f..a319cada3 100644 --- a/configs/cuda_fp16.yml +++ b/configs/cuda_fp16.yml @@ -5,6 +5,9 @@ Version: 0.0.1 Entry Point: ./numba_cuda/numba/cuda/include/12/cuda_fp16.h File List: - ./numba_cuda/numba/cuda/include/12/cuda_fp16.h +Exclude: + Struct: + - __half2 Types: __half: Number __half2: Type @@ -13,3 +16,6 @@ Data Models: __half2: StructModel Shim Include Override: "\"cuda_fp16.h\"" Require Pynvjitlink: False +Use Separate Registry: True +GPU Arch: + - sm_86 diff --git a/numba_cuda/numba/cuda/_internal/cuda_fp16.py b/numba_cuda/numba/cuda/_internal/cuda_fp16.py index 221c93510..ca7ec2a7c 100644 --- a/numba_cuda/numba/cuda/_internal/cuda_fp16.py +++ b/numba_cuda/numba/cuda/_internal/cuda_fp16.py @@ -3,13 +3,13 @@ # Automatically generated by Numbast Static Binding Generator # Generator Information: -# Ast_canopy version: 0.3.0 -# Numbast version: 0.3.0 -# Generation command: /home/lakshayg/miniconda3/envs/numbast/lib/python3.13/site-packages/numbast/__main__.py --cfg-path configs/cuda_fp16.yml --output-dir ./numba_cuda/numba/cuda/_internal/ -# Static binding generator parameters: {'cfg_path': 'configs/cuda_fp16.yml', 'output_dir': './numba_cuda/numba/cuda/_internal/', 'entry_point': None, 'retain': None, 'types': None, 'datamodels': None, 'compute_capability': None, 'run_ruff_format': True} +# Ast_canopy version: 0.5.1 +# Numbast version: 0.5.1 +# Generation command: /home/lakshayg/micromamba/envs/numbast/lib/python3.13/site-packages/numbast/__main__.py --cfg-path configs/cuda_fp16.yml --output-dir ./numba_cuda/numba/cuda/_internal/ +# Static binding generator parameters: {'cfg_path': 'configs/cuda_fp16.yml', 'output_dir': './numba_cuda/numba/cuda/_internal/', 'run_ruff_format': True} # Config file path (relative to the path of the generated binding): ../../../../../configs/cuda_fp16.yml -# Cudatoolkit version: (12, 8) -# Default CUDA_HOME path: /home/lakshayg/miniconda3/envs/numbast +# Cudatoolkit version: (12, 9) +# Default CUDA_HOME path: /home/lakshayg/micromamba/envs/numbast # Imports: @@ -23,19 +23,21 @@ from numba.core.datamodel import PrimitiveModel, StructModel from numba.core.errors import NumbaPerformanceWarning from numba.core.extending import ( + lower_cast, make_attribute_wrapper, register_model, ) -from numba.core.imputils import Registry as TargetRegistry, lower_cast -import numba.core.typeconv +from numba.core.imputils import Registry as TargetRegistry +from numba.core.imputils import lower_cast from numba.core.typing import signature +from numba.cuda import CUSource, declare_device +from numba.cuda._internal.cuda_bf16 import _type___nv_bfloat16 from numba.cuda.typing.templates import ( - Registry as TypingRegistry, AbstractTemplate, AttributeTemplate, ConcreteTemplate, ) -from numba.cuda import CUSource, declare_device +from numba.cuda.typing.templates import Registry as TypingRegistry from numba.cuda.vector_types import vector_types from numba.extending import as_numba_type from numba.types import ( @@ -58,6 +60,9 @@ ) from warnings import warn +float32x2 = vector_types["float32x2"] + + typing_registry = TypingRegistry() register = typing_registry.register register_attr = typing_registry.register_attr @@ -67,10 +72,6 @@ lower_attr = target_registry.lower_getattr lower_constant = target_registry.lower_constant -float32x2 = vector_types["float32x2"] - -# Setups: - # Shim Stream: lto_warning_raised = False @@ -108,83 +109,96 @@ def reset(self): shim_stream.write(shim_prefix) shim_obj = CUSource(shim_stream, setup_callback=lto_warning_callback) + # Enums: # Structs: -# Typing for unnamed1302257 -class _type_class_unnamed1302257(Type): +# Typing for unnamed1362071 +class _type_class_unnamed1362071(Type): def __init__(self): - super().__init__(name="unnamed1302257") + super().__init__(name="unnamed1362071") self.alignof_ = 2 self.bitwidth = 2 * 8 + def can_convert_from(self, typingctx, other): + from numba.core.typeconv import Conversion + + if other in []: + return Conversion.safe -_type_unnamed1302257 = _type_class_unnamed1302257() + +_type_unnamed1362071 = _type_class_unnamed1362071() # Make Python API for struct -unnamed1302257 = type("unnamed1302257", (), {"_nbtype": _type_unnamed1302257}) +unnamed1362071 = type("unnamed1362071", (), {"_nbtype": _type_unnamed1362071}) -as_numba_type.register(unnamed1302257, _type_unnamed1302257) +as_numba_type.register(unnamed1362071, _type_unnamed1362071) -@register_model(_type_class_unnamed1302257) -class _model_unnamed1302257(StructModel): +@register_model(_type_class_unnamed1362071) +class _model_unnamed1362071(StructModel): def __init__(self, dmm, fe_type): members = [("x", uint16)] super().__init__(dmm, fe_type, members) @register_attr -class _attr_typing_unnamed1302257(AttributeTemplate): - key = globals()["unnamed1302257"] +class _attr_typing_unnamed1362071(AttributeTemplate): + key = globals()["unnamed1362071"] def resolve_x(self, obj): return uint16 -make_attribute_wrapper(_type_class_unnamed1302257, "x", "x") +make_attribute_wrapper(_type_class_unnamed1362071, "x", "x") @register -class _ctor_template_unnamed1302257(ConcreteTemplate): - key = globals()["unnamed1302257"] +class _ctor_template_unnamed1362071(ConcreteTemplate): + key = globals()["unnamed1362071"] cases = [] -register_global(unnamed1302257, Function(_ctor_template_unnamed1302257)) +register_global(unnamed1362071, Function(_ctor_template_unnamed1362071)) -# Typing for unnamed1302366 -class _type_class_unnamed1302366(Type): +# Typing for unnamed1362180 +class _type_class_unnamed1362180(Type): def __init__(self): - super().__init__(name="unnamed1302366") + super().__init__(name="unnamed1362180") self.alignof_ = 4 self.bitwidth = 4 * 8 + def can_convert_from(self, typingctx, other): + from numba.core.typeconv import Conversion + + if other in []: + return Conversion.safe -_type_unnamed1302366 = _type_class_unnamed1302366() + +_type_unnamed1362180 = _type_class_unnamed1362180() # Make Python API for struct -unnamed1302366 = type("unnamed1302366", (), {"_nbtype": _type_unnamed1302366}) +unnamed1362180 = type("unnamed1362180", (), {"_nbtype": _type_unnamed1362180}) -as_numba_type.register(unnamed1302366, _type_unnamed1302366) +as_numba_type.register(unnamed1362180, _type_unnamed1362180) -@register_model(_type_class_unnamed1302366) -class _model_unnamed1302366(StructModel): +@register_model(_type_class_unnamed1362180) +class _model_unnamed1362180(StructModel): def __init__(self, dmm, fe_type): members = [("x", uint16), ("y", uint16)] super().__init__(dmm, fe_type, members) @register_attr -class _attr_typing_unnamed1302366(AttributeTemplate): - key = globals()["unnamed1302366"] +class _attr_typing_unnamed1362180(AttributeTemplate): + key = globals()["unnamed1362180"] def resolve_x(self, obj): return uint16 @@ -193,20 +207,19 @@ def resolve_y(self, obj): return uint16 -make_attribute_wrapper(_type_class_unnamed1302366, "x", "x") +make_attribute_wrapper(_type_class_unnamed1362180, "x", "x") -make_attribute_wrapper(_type_class_unnamed1302366, "y", "y") +make_attribute_wrapper(_type_class_unnamed1362180, "y", "y") @register -class _ctor_template_unnamed1302366(ConcreteTemplate): - key = globals()["unnamed1302366"] +class _ctor_template_unnamed1362180(ConcreteTemplate): + key = globals()["unnamed1362180"] cases = [] -register_global(unnamed1302366, Function(_ctor_template_unnamed1302366)) - +register_global(unnamed1362180, Function(_ctor_template_unnamed1362180)) __half = _type___half = numba.core.types.float16 setattr(__half, "alignof_", 2) @@ -216,14 +229,14 @@ class _ctor_template_unnamed1302366(ConcreteTemplate): def _lower__ZN6__halfC1Ev(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Ev_1(int &ignore, __half *self ) { + _ZN6__halfC1Ev_nbst(int &ignore, __half *self ) { new (self) __half(); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Ev_1", + "_ZN6__halfC1Ev_nbst", int32( CPointer(_type___half), ), @@ -237,7 +250,7 @@ def __half_device_caller(arg_0): ) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Ev_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Ev_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -267,25 +280,25 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1ERK10__half_raw(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1ERK10__half_raw_1(int &ignore, __half *self , __half_raw* hr) { + _ZN6__halfC1ERK10__half_raw_nbst(int &ignore, __half *self , __half_raw* hr) { new (self) __half(*hr); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1ERK10__half_raw_1", - int32(CPointer(_type___half), CPointer(_type_unnamed1302257)), + "_ZN6__halfC1ERK10__half_raw_nbst", + int32(CPointer(_type___half), CPointer(_type_unnamed1362071)), ) def __half_device_caller(arg_0, arg_1): return _ctor_decl___half(arg_0, arg_1) - @lower(__half, _type_unnamed1302257) + @lower(__half, _type_unnamed1362071) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) shim_stream.write_with_key( - "____half___ZN6__halfC1ERK10__half_raw_1", shim_raw_str + "_ZN6__halfC1ERK10__half_raw_nbst", shim_raw_str ) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" @@ -300,7 +313,7 @@ def ctor_impl(context, builder, sig, args): builder, __half_device_caller, signature( - int32, CPointer(_type___half), CPointer(_type_unnamed1302257) + int32, CPointer(_type___half), CPointer(_type_unnamed1362071) ), (selfptr, *argptrs), ) @@ -312,18 +325,65 @@ def ctor_impl(context, builder, sig, args): _lower__ZN6__halfC1ERK10__half_raw(shim_stream, shim_obj) +def _lower__ZN6__halfC1E13__nv_bfloat16(shim_stream, shim_obj): + shim_raw_str = """ + extern "C" __device__ int + _ZN6__halfC1E13__nv_bfloat16_nbst(int &ignore, __half *self , __nv_bfloat16* f) { + new (self) __half(*f); + return 0; + } + """ + + _ctor_decl___half = declare_device( + "_ZN6__halfC1E13__nv_bfloat16_nbst", + int32(CPointer(_type___half), CPointer(_type___nv_bfloat16)), + ) + + def __half_device_caller(arg_0, arg_1): + return _ctor_decl___half(arg_0, arg_1) + + @lower(__half, _type___nv_bfloat16) + def ctor_impl(context, builder, sig, args): + context.active_code_library.add_linking_file(shim_obj) + shim_stream.write_with_key( + "_ZN6__halfC1E13__nv_bfloat16_nbst", shim_raw_str + ) + selfptr = builder.alloca( + context.get_value_type(_type___half), name="selfptr" + ) + argptrs = [ + builder.alloca(context.get_value_type(arg)) for arg in sig.args + ] + for ptr, ty, arg in zip(argptrs, sig.args, args): + builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) + + context.compile_internal( + builder, + __half_device_caller, + signature( + int32, CPointer(_type___half), CPointer(_type___nv_bfloat16) + ), + (selfptr, *argptrs), + ) + return builder.load( + selfptr, align=getattr(_type___half, "alignof_", None) + ) + + +_lower__ZN6__halfC1E13__nv_bfloat16(shim_stream, shim_obj) + + def _lower__ZN6__halfC1Ef(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Ef_1(int &ignore, __half *self , float* f) { + _ZN6__halfC1Ef_nbst(int &ignore, __half *self , float* f) { new (self) __half(*f); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Ef_1", - int32(CPointer(_type___half), CPointer(float32)), + "_ZN6__halfC1Ef_nbst", int32(CPointer(_type___half), CPointer(float32)) ) def __half_device_caller(arg_0, arg_1): @@ -332,7 +392,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, float32) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Ef_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Ef_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -359,15 +419,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Ed(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Ed_1(int &ignore, __half *self , double* f) { + _ZN6__halfC1Ed_nbst(int &ignore, __half *self , double* f) { new (self) __half(*f); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Ed_1", - int32(CPointer(_type___half), CPointer(float64)), + "_ZN6__halfC1Ed_nbst", int32(CPointer(_type___half), CPointer(float64)) ) def __half_device_caller(arg_0, arg_1): @@ -376,7 +435,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, float64) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Ed_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Ed_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -403,15 +462,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Es(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Es_1(int &ignore, __half *self , short* val) { + _ZN6__halfC1Es_nbst(int &ignore, __half *self , short* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Es_1", - int32(CPointer(_type___half), CPointer(int16)), + "_ZN6__halfC1Es_nbst", int32(CPointer(_type___half), CPointer(int16)) ) def __half_device_caller(arg_0, arg_1): @@ -420,7 +478,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, int16) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Es_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Es_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -447,15 +505,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Et(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Et_1(int &ignore, __half *self , unsigned short* val) { + _ZN6__halfC1Et_nbst(int &ignore, __half *self , unsigned short* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Et_1", - int32(CPointer(_type___half), CPointer(uint16)), + "_ZN6__halfC1Et_nbst", int32(CPointer(_type___half), CPointer(uint16)) ) def __half_device_caller(arg_0, arg_1): @@ -464,7 +521,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, uint16) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Et_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Et_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -491,15 +548,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Ei(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Ei_1(int &ignore, __half *self , int* val) { + _ZN6__halfC1Ei_nbst(int &ignore, __half *self , int* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Ei_1", - int32(CPointer(_type___half), CPointer(int32)), + "_ZN6__halfC1Ei_nbst", int32(CPointer(_type___half), CPointer(int32)) ) def __half_device_caller(arg_0, arg_1): @@ -508,7 +564,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, int32) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Ei_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Ei_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -535,15 +591,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Ej(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Ej_1(int &ignore, __half *self , unsigned int* val) { + _ZN6__halfC1Ej_nbst(int &ignore, __half *self , unsigned int* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Ej_1", - int32(CPointer(_type___half), CPointer(uint32)), + "_ZN6__halfC1Ej_nbst", int32(CPointer(_type___half), CPointer(uint32)) ) def __half_device_caller(arg_0, arg_1): @@ -552,7 +607,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, uint32) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Ej_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Ej_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -579,15 +634,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1El(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1El_1(int &ignore, __half *self , long* val) { + _ZN6__halfC1El_nbst(int &ignore, __half *self , long* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1El_1", - int32(CPointer(_type___half), CPointer(int64)), + "_ZN6__halfC1El_nbst", int32(CPointer(_type___half), CPointer(int64)) ) def __half_device_caller(arg_0, arg_1): @@ -596,7 +650,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, int64) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1El_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1El_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -623,15 +677,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Em(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Em_1(int &ignore, __half *self , unsigned long* val) { + _ZN6__halfC1Em_nbst(int &ignore, __half *self , unsigned long* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Em_1", - int32(CPointer(_type___half), CPointer(uint64)), + "_ZN6__halfC1Em_nbst", int32(CPointer(_type___half), CPointer(uint64)) ) def __half_device_caller(arg_0, arg_1): @@ -640,7 +693,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, uint64) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Em_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Em_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -667,15 +720,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Ex(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Ex_1(int &ignore, __half *self , long long* val) { + _ZN6__halfC1Ex_nbst(int &ignore, __half *self , long long* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Ex_1", - int32(CPointer(_type___half), CPointer(int64)), + "_ZN6__halfC1Ex_nbst", int32(CPointer(_type___half), CPointer(int64)) ) def __half_device_caller(arg_0, arg_1): @@ -684,7 +736,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, int64) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Ex_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Ex_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -711,15 +763,14 @@ def ctor_impl(context, builder, sig, args): def _lower__ZN6__halfC1Ey(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half___ZN6__halfC1Ey_1(int &ignore, __half *self , unsigned long long* val) { + _ZN6__halfC1Ey_nbst(int &ignore, __half *self , unsigned long long* val) { new (self) __half(*val); return 0; } """ _ctor_decl___half = declare_device( - "____half___ZN6__halfC1Ey_1", - int32(CPointer(_type___half), CPointer(uint64)), + "_ZN6__halfC1Ey_nbst", int32(CPointer(_type___half), CPointer(uint64)) ) def __half_device_caller(arg_0, arg_1): @@ -728,7 +779,7 @@ def __half_device_caller(arg_0, arg_1): @lower(__half, uint64) def ctor_impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half___ZN6__halfC1Ey_1", shim_raw_str) + shim_stream.write_with_key("_ZN6__halfC1Ey_nbst", shim_raw_str) selfptr = builder.alloca( context.get_value_type(_type___half), name="selfptr" ) @@ -759,7 +810,8 @@ class _ctor_template___half(ConcreteTemplate): signature( _type___half, ), - signature(_type___half, _type_unnamed1302257), + signature(_type___half, _type_unnamed1362071), + signature(_type___half, _type___nv_bfloat16), signature(_type___half, float32), signature(_type___half, float64), signature(_type___half, int16), @@ -776,7 +828,7 @@ class _ctor_template___half(ConcreteTemplate): register_global(__half, Function(_ctor_template___half)) -def _from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj): +def _from___half_to__type_unnamed1362071_lower(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int ____half__ZNK6__halfcv10__half_rawEv_1(__half_raw &retval, __half *self) { @@ -787,7 +839,7 @@ def _from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj): _op_decl___half = declare_device( "____half__ZNK6__halfcv10__half_rawEv_1", - _type_unnamed1302257( + _type_unnamed1362071( CPointer(_type___half), ), ) @@ -795,7 +847,7 @@ def _from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj): def _conversion_op_caller___half(arg): return _op_decl___half(arg) - @lower_cast(_type___half, _type_unnamed1302257) + @lower_cast(_type___half, _type_unnamed1362071) def impl(context, builder, fromty, toty, value): context.active_code_library.add_linking_file(shim_obj) shim_stream.write_with_key( @@ -810,17 +862,17 @@ def impl(context, builder, fromty, toty, value): builder, _conversion_op_caller___half, signature( - _type_unnamed1302257, + _type_unnamed1362071, CPointer(_type___half), ), (ptr,), ) -_from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj) +_from___half_to__type_unnamed1362071_lower(shim_stream, shim_obj) -def _from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj): +def _from___half_to__type_unnamed1362071_lower(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int ____half__ZNVK6__halfcv10__half_rawEv_1(__half_raw &retval, __half *self) { @@ -831,7 +883,7 @@ def _from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj): _op_decl___half = declare_device( "____half__ZNVK6__halfcv10__half_rawEv_1", - _type_unnamed1302257( + _type_unnamed1362071( CPointer(_type___half), ), ) @@ -839,7 +891,7 @@ def _from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj): def _conversion_op_caller___half(arg): return _op_decl___half(arg) - @lower_cast(_type___half, _type_unnamed1302257) + @lower_cast(_type___half, _type_unnamed1362071) def impl(context, builder, fromty, toty, value): context.active_code_library.add_linking_file(shim_obj) shim_stream.write_with_key( @@ -854,1578 +906,1484 @@ def impl(context, builder, fromty, toty, value): builder, _conversion_op_caller___half, signature( - _type_unnamed1302257, + _type_unnamed1362071, CPointer(_type___half), ), (ptr,), ) -_from___half_to__type_unnamed1302257_lower(shim_stream, shim_obj) - - -# Typing for __half2 -class _type_class___half2(Type): - def __init__(self): - super().__init__(name="__half2") - self.alignof_ = 4 - self.bitwidth = 4 * 8 - - -_type___half2 = _type_class___half2() - - -# Make Python API for struct -__half2 = type("__half2", (), {"_nbtype": _type___half2}) - -as_numba_type.register(__half2, _type___half2) - - -@register_model(_type_class___half2) -class _model___half2(StructModel): - def __init__(self, dmm, fe_type): - members = [("x", _type___half), ("y", _type___half)] - super().__init__(dmm, fe_type, members) - - -@register_attr -class _attr_typing___half2(AttributeTemplate): - key = globals()["__half2"] - - def resolve_x(self, obj): - return _type___half - - def resolve_y(self, obj): - return _type___half +_from___half_to__type_unnamed1362071_lower(shim_stream, shim_obj) -make_attribute_wrapper(_type_class___half2, "x", "x") +# Functions: -make_attribute_wrapper(_type_class___half2, "y", "y") +def __double2half(): + pass -def _lower__ZN7__half2C1Ev(shim_stream, shim_obj): +def _lower__ZL13__double2halfd_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half2___ZN7__half2C1Ev_1(int &ignore, __half2 *self ) { - new (self) __half2(); + _ZL13__double2halfd_nbst(__half &retval , double* a) { + retval = __double2half(*a); return 0; } """ - _ctor_decl___half2 = declare_device( - "____half2___ZN7__half2C1Ev_1", - int32( - CPointer(_type___half2), - ), + _ZL13__double2halfd_nbst = declare_device( + "_ZL13__double2halfd_nbst", _type___half(CPointer(float64)) ) - def __half2_device_caller(arg_0): - return _ctor_decl___half2(arg_0) + def _ZL13__double2halfd_nbst_caller(arg_0): + return _ZL13__double2halfd_nbst(arg_0) - @lower( - __half2, - ) - def ctor_impl(context, builder, sig, args): + @lower(__double2half, float64) + def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("____half2___ZN7__half2C1Ev_1", shim_raw_str) - selfptr = builder.alloca( - context.get_value_type(_type___half2), name="selfptr" - ) - argptrs = [ - builder.alloca(context.get_value_type(arg)) for arg in sig.args - ] - for ptr, ty, arg in zip(argptrs, sig.args, args): + shim_stream.write_with_key("_ZL13__double2halfd_nbst", shim_raw_str) + ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] + for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - context.compile_internal( + return context.compile_internal( builder, - __half2_device_caller, - signature( - int32, - CPointer(_type___half2), - ), - (selfptr, *argptrs), - ) - return builder.load( - selfptr, align=getattr(_type___half2, "alignof_", None) + _ZL13__double2halfd_nbst_caller, + signature(_type___half, CPointer(float64)), + ptrs, ) -_lower__ZN7__half2C1Ev(shim_stream, shim_obj) +_lower__ZL13__double2halfd_nbst(shim_stream, shim_obj) + + +def __float2half(): + pass -def _lower__ZN7__half2C1EOKS_(shim_stream, shim_obj): +def _lower__ZL12__float2halff_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half2___ZN7__half2C1EOKS__1(int &ignore, __half2 *self , __half2* src) { - new (self) __half2(*src); + _ZL12__float2halff_nbst(__half &retval , float* a) { + retval = __float2half(*a); return 0; } """ - _ctor_decl___half2 = declare_device( - "____half2___ZN7__half2C1EOKS__1", - int32(CPointer(_type___half2), CPointer(_type___half2)), + _ZL12__float2halff_nbst = declare_device( + "_ZL12__float2halff_nbst", _type___half(CPointer(float32)) ) - def __half2_device_caller(arg_0, arg_1): - return _ctor_decl___half2(arg_0, arg_1) + def _ZL12__float2halff_nbst_caller(arg_0): + return _ZL12__float2halff_nbst(arg_0) - @lower(__half2, _type___half2) - def ctor_impl(context, builder, sig, args): + @lower(__float2half, float32) + def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "____half2___ZN7__half2C1EOKS__1", shim_raw_str - ) - selfptr = builder.alloca( - context.get_value_type(_type___half2), name="selfptr" - ) - argptrs = [ - builder.alloca(context.get_value_type(arg)) for arg in sig.args - ] - for ptr, ty, arg in zip(argptrs, sig.args, args): + shim_stream.write_with_key("_ZL12__float2halff_nbst", shim_raw_str) + ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] + for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - context.compile_internal( + return context.compile_internal( builder, - __half2_device_caller, - signature(int32, CPointer(_type___half2), CPointer(_type___half2)), - (selfptr, *argptrs), - ) - return builder.load( - selfptr, align=getattr(_type___half2, "alignof_", None) + _ZL12__float2halff_nbst_caller, + signature(_type___half, CPointer(float32)), + ptrs, ) -_lower__ZN7__half2C1EOKS_(shim_stream, shim_obj) +_lower__ZL12__float2halff_nbst(shim_stream, shim_obj) + + +def __float2half_rn(): + pass -def _lower__ZN7__half2C1ERK6__halfS2_(shim_stream, shim_obj): +def _lower__ZL15__float2half_rnf_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half2___ZN7__half2C1ERK6__halfS2__1(int &ignore, __half2 *self , __half* a, __half* b) { - new (self) __half2(*a, *b); + _ZL15__float2half_rnf_nbst(__half &retval , float* a) { + retval = __float2half_rn(*a); return 0; } """ - _ctor_decl___half2 = declare_device( - "____half2___ZN7__half2C1ERK6__halfS2__1", - int32( - CPointer(_type___half2), - CPointer(_type___half), - CPointer(_type___half), - ), + _ZL15__float2half_rnf_nbst = declare_device( + "_ZL15__float2half_rnf_nbst", _type___half(CPointer(float32)) ) - def __half2_device_caller(arg_0, arg_1, arg_2): - return _ctor_decl___half2(arg_0, arg_1, arg_2) + def _ZL15__float2half_rnf_nbst_caller(arg_0): + return _ZL15__float2half_rnf_nbst(arg_0) - @lower(__half2, _type___half, _type___half) - def ctor_impl(context, builder, sig, args): + @lower(__float2half_rn, float32) + def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "____half2___ZN7__half2C1ERK6__halfS2__1", shim_raw_str - ) - selfptr = builder.alloca( - context.get_value_type(_type___half2), name="selfptr" - ) - argptrs = [ - builder.alloca(context.get_value_type(arg)) for arg in sig.args - ] - for ptr, ty, arg in zip(argptrs, sig.args, args): + shim_stream.write_with_key("_ZL15__float2half_rnf_nbst", shim_raw_str) + ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] + for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - context.compile_internal( + return context.compile_internal( builder, - __half2_device_caller, - signature( - int32, - CPointer(_type___half2), - CPointer(_type___half), - CPointer(_type___half), - ), - (selfptr, *argptrs), - ) - return builder.load( - selfptr, align=getattr(_type___half2, "alignof_", None) + _ZL15__float2half_rnf_nbst_caller, + signature(_type___half, CPointer(float32)), + ptrs, ) -_lower__ZN7__half2C1ERK6__halfS2_(shim_stream, shim_obj) +_lower__ZL15__float2half_rnf_nbst(shim_stream, shim_obj) + + +def __float2half_rz(): + pass -def _lower__ZN7__half2C1ERKS_(shim_stream, shim_obj): +def _lower__ZL15__float2half_rzf_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half2___ZN7__half2C1ERKS__1(int &ignore, __half2 *self , __half2* src) { - new (self) __half2(*src); + _ZL15__float2half_rzf_nbst(__half &retval , float* a) { + retval = __float2half_rz(*a); return 0; } """ - _ctor_decl___half2 = declare_device( - "____half2___ZN7__half2C1ERKS__1", - int32(CPointer(_type___half2), CPointer(_type___half2)), + _ZL15__float2half_rzf_nbst = declare_device( + "_ZL15__float2half_rzf_nbst", _type___half(CPointer(float32)) ) - def __half2_device_caller(arg_0, arg_1): - return _ctor_decl___half2(arg_0, arg_1) + def _ZL15__float2half_rzf_nbst_caller(arg_0): + return _ZL15__float2half_rzf_nbst(arg_0) - @lower(__half2, _type___half2) - def ctor_impl(context, builder, sig, args): + @lower(__float2half_rz, float32) + def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "____half2___ZN7__half2C1ERKS__1", shim_raw_str - ) - selfptr = builder.alloca( - context.get_value_type(_type___half2), name="selfptr" - ) - argptrs = [ - builder.alloca(context.get_value_type(arg)) for arg in sig.args - ] - for ptr, ty, arg in zip(argptrs, sig.args, args): + shim_stream.write_with_key("_ZL15__float2half_rzf_nbst", shim_raw_str) + ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] + for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - context.compile_internal( + return context.compile_internal( builder, - __half2_device_caller, - signature(int32, CPointer(_type___half2), CPointer(_type___half2)), - (selfptr, *argptrs), - ) - return builder.load( - selfptr, align=getattr(_type___half2, "alignof_", None) + _ZL15__float2half_rzf_nbst_caller, + signature(_type___half, CPointer(float32)), + ptrs, ) -_lower__ZN7__half2C1ERKS_(shim_stream, shim_obj) +_lower__ZL15__float2half_rzf_nbst(shim_stream, shim_obj) + + +def __float2half_rd(): + pass -def _lower__ZN7__half2C1ERK11__half2_raw(shim_stream, shim_obj): +def _lower__ZL15__float2half_rdf_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half2___ZN7__half2C1ERK11__half2_raw_1(int &ignore, __half2 *self , __half2_raw* h2r) { - new (self) __half2(*h2r); + _ZL15__float2half_rdf_nbst(__half &retval , float* a) { + retval = __float2half_rd(*a); return 0; } """ - _ctor_decl___half2 = declare_device( - "____half2___ZN7__half2C1ERK11__half2_raw_1", - int32(CPointer(_type___half2), CPointer(_type_unnamed1302366)), + _ZL15__float2half_rdf_nbst = declare_device( + "_ZL15__float2half_rdf_nbst", _type___half(CPointer(float32)) ) - def __half2_device_caller(arg_0, arg_1): - return _ctor_decl___half2(arg_0, arg_1) + def _ZL15__float2half_rdf_nbst_caller(arg_0): + return _ZL15__float2half_rdf_nbst(arg_0) - @lower(__half2, _type_unnamed1302366) - def ctor_impl(context, builder, sig, args): + @lower(__float2half_rd, float32) + def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "____half2___ZN7__half2C1ERK11__half2_raw_1", shim_raw_str - ) - selfptr = builder.alloca( - context.get_value_type(_type___half2), name="selfptr" - ) - argptrs = [ - builder.alloca(context.get_value_type(arg)) for arg in sig.args - ] - for ptr, ty, arg in zip(argptrs, sig.args, args): + shim_stream.write_with_key("_ZL15__float2half_rdf_nbst", shim_raw_str) + ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] + for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - context.compile_internal( + return context.compile_internal( builder, - __half2_device_caller, - signature( - int32, CPointer(_type___half2), CPointer(_type_unnamed1302366) - ), - (selfptr, *argptrs), - ) - return builder.load( - selfptr, align=getattr(_type___half2, "alignof_", None) + _ZL15__float2half_rdf_nbst_caller, + signature(_type___half, CPointer(float32)), + ptrs, ) -_lower__ZN7__half2C1ERK11__half2_raw(shim_stream, shim_obj) +_lower__ZL15__float2half_rdf_nbst(shim_stream, shim_obj) -@register -class _ctor_template___half2(ConcreteTemplate): - key = globals()["__half2"] - cases = [ - signature( - _type___half2, - ), - signature(_type___half2, _type___half2), - signature(_type___half2, _type___half, _type___half), - signature(_type___half2, _type___half2), - signature(_type___half2, _type_unnamed1302366), - ] - - -register_global(__half2, Function(_ctor_template___half2)) +def __float2half_ru(): + pass -def _from___half2_to__type_unnamed1302366_lower(shim_stream, shim_obj): +def _lower__ZL15__float2half_ruf_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - ____half2__ZNK7__half2cv11__half2_rawEv_1(__half2_raw &retval, __half2 *self) { - retval = self->operator __half2_raw(); + _ZL15__float2half_ruf_nbst(__half &retval , float* a) { + retval = __float2half_ru(*a); return 0; } """ - _op_decl___half2 = declare_device( - "____half2__ZNK7__half2cv11__half2_rawEv_1", - _type_unnamed1302366( - CPointer(_type___half2), - ), + _ZL15__float2half_ruf_nbst = declare_device( + "_ZL15__float2half_ruf_nbst", _type___half(CPointer(float32)) ) - def _conversion_op_caller___half2(arg): - return _op_decl___half2(arg) + def _ZL15__float2half_ruf_nbst_caller(arg_0): + return _ZL15__float2half_ruf_nbst(arg_0) - @lower_cast(_type___half2, _type_unnamed1302366) - def impl(context, builder, fromty, toty, value): + @lower(__float2half_ru, float32) + def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "____half2__ZNK7__half2cv11__half2_rawEv_1", shim_raw_str - ) - ptr = builder.alloca( - context.get_value_type(_type___half2), name="selfptr" - ) - builder.store(value, ptr, align=getattr(_type___half2, "align", None)) + shim_stream.write_with_key("_ZL15__float2half_ruf_nbst", shim_raw_str) + ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] + for ptr, ty, arg in zip(ptrs, sig.args, args): + builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _conversion_op_caller___half2, - signature( - _type_unnamed1302366, - CPointer(_type___half2), - ), - (ptr,), + _ZL15__float2half_ruf_nbst_caller, + signature(_type___half, CPointer(float32)), + ptrs, ) -_from___half2_to__type_unnamed1302366_lower(shim_stream, shim_obj) - - -# Functions: +_lower__ZL15__float2half_ruf_nbst(shim_stream, shim_obj) -def __double2half(): +def __half2float(): pass -def _lower__ZL13__double2halfd_1(shim_stream, shim_obj): +def _lower__ZL12__half2float6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__double2halfd_1(__half &retval , double* a) { - retval = __double2half(*a); + _ZL12__half2float6__half_nbst(float &retval , __half* a) { + retval = __half2float(*a); return 0; } """ - _ZL13__double2halfd_1 = declare_device( - "_ZL13__double2halfd_1", _type___half(CPointer(float64)) + _ZL12__half2float6__half_nbst = declare_device( + "_ZL12__half2float6__half_nbst", float32(CPointer(_type___half)) ) - def _ZL13__double2halfd_1_caller(arg_0): - return _ZL13__double2halfd_1(arg_0) + def _ZL12__half2float6__half_nbst_caller(arg_0): + return _ZL12__half2float6__half_nbst(arg_0) - @lower(__double2half, float64) + @lower(__half2float, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__double2halfd_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL12__half2float6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__double2halfd_1_caller, - signature(_type___half, CPointer(float64)), + _ZL12__half2float6__half_nbst_caller, + signature(float32, CPointer(_type___half)), ptrs, ) -_lower__ZL13__double2halfd_1(shim_stream, shim_obj) +_lower__ZL12__half2float6__half_nbst(shim_stream, shim_obj) -def __float2half(): +def __half2char_rz(): pass -def _lower__ZL12__float2halff_1(shim_stream, shim_obj): +def _lower__ZL14__half2char_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__float2halff_1(__half &retval , float* a) { - retval = __float2half(*a); + _ZL14__half2char_rz6__half_nbst(signed char &retval , __half* h) { + retval = __half2char_rz(*h); return 0; } """ - _ZL12__float2halff_1 = declare_device( - "_ZL12__float2halff_1", _type___half(CPointer(float32)) + _ZL14__half2char_rz6__half_nbst = declare_device( + "_ZL14__half2char_rz6__half_nbst", int8(CPointer(_type___half)) ) - def _ZL12__float2halff_1_caller(arg_0): - return _ZL12__float2halff_1(arg_0) + def _ZL14__half2char_rz6__half_nbst_caller(arg_0): + return _ZL14__half2char_rz6__half_nbst(arg_0) - @lower(__float2half, float32) + @lower(__half2char_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__float2halff_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL14__half2char_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__float2halff_1_caller, - signature(_type___half, CPointer(float32)), + _ZL14__half2char_rz6__half_nbst_caller, + signature(int8, CPointer(_type___half)), ptrs, ) -_lower__ZL12__float2halff_1(shim_stream, shim_obj) +_lower__ZL14__half2char_rz6__half_nbst(shim_stream, shim_obj) -def __float2half_rn(): +def __half2uchar_rz(): pass -def _lower__ZL15__float2half_rnf_1(shim_stream, shim_obj): +def _lower__ZL15__half2uchar_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__float2half_rnf_1(__half &retval , float* a) { - retval = __float2half_rn(*a); + _ZL15__half2uchar_rz6__half_nbst(unsigned char &retval , __half* h) { + retval = __half2uchar_rz(*h); return 0; } """ - _ZL15__float2half_rnf_1 = declare_device( - "_ZL15__float2half_rnf_1", _type___half(CPointer(float32)) + _ZL15__half2uchar_rz6__half_nbst = declare_device( + "_ZL15__half2uchar_rz6__half_nbst", uint8(CPointer(_type___half)) ) - def _ZL15__float2half_rnf_1_caller(arg_0): - return _ZL15__float2half_rnf_1(arg_0) + def _ZL15__half2uchar_rz6__half_nbst_caller(arg_0): + return _ZL15__half2uchar_rz6__half_nbst(arg_0) - @lower(__float2half_rn, float32) + @lower(__half2uchar_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__float2half_rnf_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL15__half2uchar_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__float2half_rnf_1_caller, - signature(_type___half, CPointer(float32)), + _ZL15__half2uchar_rz6__half_nbst_caller, + signature(uint8, CPointer(_type___half)), ptrs, ) -_lower__ZL15__float2half_rnf_1(shim_stream, shim_obj) +_lower__ZL15__half2uchar_rz6__half_nbst(shim_stream, shim_obj) -def __float2half_rz(): +def __half2short_rz(): pass -def _lower__ZL15__float2half_rzf_1(shim_stream, shim_obj): +def _lower__ZL15__half2short_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__float2half_rzf_1(__half &retval , float* a) { - retval = __float2half_rz(*a); + _ZL15__half2short_rz6__half_nbst(short &retval , __half* h) { + retval = __half2short_rz(*h); return 0; } """ - _ZL15__float2half_rzf_1 = declare_device( - "_ZL15__float2half_rzf_1", _type___half(CPointer(float32)) + _ZL15__half2short_rz6__half_nbst = declare_device( + "_ZL15__half2short_rz6__half_nbst", int16(CPointer(_type___half)) ) - def _ZL15__float2half_rzf_1_caller(arg_0): - return _ZL15__float2half_rzf_1(arg_0) + def _ZL15__half2short_rz6__half_nbst_caller(arg_0): + return _ZL15__half2short_rz6__half_nbst(arg_0) - @lower(__float2half_rz, float32) + @lower(__half2short_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__float2half_rzf_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL15__half2short_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__float2half_rzf_1_caller, - signature(_type___half, CPointer(float32)), + _ZL15__half2short_rz6__half_nbst_caller, + signature(int16, CPointer(_type___half)), ptrs, ) -_lower__ZL15__float2half_rzf_1(shim_stream, shim_obj) +_lower__ZL15__half2short_rz6__half_nbst(shim_stream, shim_obj) -def __float2half_rd(): +def __half2ushort_rz(): pass -def _lower__ZL15__float2half_rdf_1(shim_stream, shim_obj): +def _lower__ZL16__half2ushort_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__float2half_rdf_1(__half &retval , float* a) { - retval = __float2half_rd(*a); + _ZL16__half2ushort_rz6__half_nbst(unsigned short &retval , __half* h) { + retval = __half2ushort_rz(*h); return 0; } """ - _ZL15__float2half_rdf_1 = declare_device( - "_ZL15__float2half_rdf_1", _type___half(CPointer(float32)) + _ZL16__half2ushort_rz6__half_nbst = declare_device( + "_ZL16__half2ushort_rz6__half_nbst", uint16(CPointer(_type___half)) ) - def _ZL15__float2half_rdf_1_caller(arg_0): - return _ZL15__float2half_rdf_1(arg_0) + def _ZL16__half2ushort_rz6__half_nbst_caller(arg_0): + return _ZL16__half2ushort_rz6__half_nbst(arg_0) - @lower(__float2half_rd, float32) + @lower(__half2ushort_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__float2half_rdf_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL16__half2ushort_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__float2half_rdf_1_caller, - signature(_type___half, CPointer(float32)), + _ZL16__half2ushort_rz6__half_nbst_caller, + signature(uint16, CPointer(_type___half)), ptrs, ) -_lower__ZL15__float2half_rdf_1(shim_stream, shim_obj) +_lower__ZL16__half2ushort_rz6__half_nbst(shim_stream, shim_obj) -def __float2half_ru(): +def __half2int_rz(): pass -def _lower__ZL15__float2half_ruf_1(shim_stream, shim_obj): +def _lower__ZL13__half2int_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__float2half_ruf_1(__half &retval , float* a) { - retval = __float2half_ru(*a); + _ZL13__half2int_rz6__half_nbst(int &retval , __half* h) { + retval = __half2int_rz(*h); return 0; } """ - _ZL15__float2half_ruf_1 = declare_device( - "_ZL15__float2half_ruf_1", _type___half(CPointer(float32)) + _ZL13__half2int_rz6__half_nbst = declare_device( + "_ZL13__half2int_rz6__half_nbst", int32(CPointer(_type___half)) ) - def _ZL15__float2half_ruf_1_caller(arg_0): - return _ZL15__float2half_ruf_1(arg_0) + def _ZL13__half2int_rz6__half_nbst_caller(arg_0): + return _ZL13__half2int_rz6__half_nbst(arg_0) - @lower(__float2half_ru, float32) + @lower(__half2int_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__float2half_ruf_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2int_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__float2half_ruf_1_caller, - signature(_type___half, CPointer(float32)), + _ZL13__half2int_rz6__half_nbst_caller, + signature(int32, CPointer(_type___half)), ptrs, ) -_lower__ZL15__float2half_ruf_1(shim_stream, shim_obj) +_lower__ZL13__half2int_rz6__half_nbst(shim_stream, shim_obj) -def __half2float(): +def __half2uint_rz(): pass -def _lower__ZL12__half2float6__half_1(shim_stream, shim_obj): +def _lower__ZL14__half2uint_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__half2float6__half_1(float &retval , __half* a) { - retval = __half2float(*a); + _ZL14__half2uint_rz6__half_nbst(unsigned int &retval , __half* h) { + retval = __half2uint_rz(*h); return 0; } """ - _ZL12__half2float6__half_1 = declare_device( - "_ZL12__half2float6__half_1", float32(CPointer(_type___half)) + _ZL14__half2uint_rz6__half_nbst = declare_device( + "_ZL14__half2uint_rz6__half_nbst", uint32(CPointer(_type___half)) ) - def _ZL12__half2float6__half_1_caller(arg_0): - return _ZL12__half2float6__half_1(arg_0) + def _ZL14__half2uint_rz6__half_nbst_caller(arg_0): + return _ZL14__half2uint_rz6__half_nbst(arg_0) - @lower(__half2float, _type___half) + @lower(__half2uint_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__half2float6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL14__half2uint_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__half2float6__half_1_caller, - signature(float32, CPointer(_type___half)), + _ZL14__half2uint_rz6__half_nbst_caller, + signature(uint32, CPointer(_type___half)), ptrs, ) -_lower__ZL12__half2float6__half_1(shim_stream, shim_obj) +_lower__ZL14__half2uint_rz6__half_nbst(shim_stream, shim_obj) -def __float2half2_rn(): +def __half2ll_rz(): pass -def _lower__ZL16__float2half2_rnf_1(shim_stream, shim_obj): +def _lower__ZL12__half2ll_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__float2half2_rnf_1(__half2 &retval , float* a) { - retval = __float2half2_rn(*a); + _ZL12__half2ll_rz6__half_nbst(long long &retval , __half* h) { + retval = __half2ll_rz(*h); return 0; } """ - _ZL16__float2half2_rnf_1 = declare_device( - "_ZL16__float2half2_rnf_1", _type___half2(CPointer(float32)) + _ZL12__half2ll_rz6__half_nbst = declare_device( + "_ZL12__half2ll_rz6__half_nbst", int64(CPointer(_type___half)) ) - def _ZL16__float2half2_rnf_1_caller(arg_0): - return _ZL16__float2half2_rnf_1(arg_0) + def _ZL12__half2ll_rz6__half_nbst_caller(arg_0): + return _ZL12__half2ll_rz6__half_nbst(arg_0) - @lower(__float2half2_rn, float32) + @lower(__half2ll_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL16__float2half2_rnf_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL12__half2ll_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__float2half2_rnf_1_caller, - signature(_type___half2, CPointer(float32)), + _ZL12__half2ll_rz6__half_nbst_caller, + signature(int64, CPointer(_type___half)), ptrs, ) -_lower__ZL16__float2half2_rnf_1(shim_stream, shim_obj) +_lower__ZL12__half2ll_rz6__half_nbst(shim_stream, shim_obj) -def __floats2half2_rn(): +def __half2ull_rz(): pass -def _lower__ZL17__floats2half2_rnff_1(shim_stream, shim_obj): +def _lower__ZL13__half2ull_rz6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL17__floats2half2_rnff_1(__half2 &retval , float* a, float* b) { - retval = __floats2half2_rn(*a, *b); + _ZL13__half2ull_rz6__half_nbst(unsigned long long &retval , __half* h) { + retval = __half2ull_rz(*h); return 0; } """ - _ZL17__floats2half2_rnff_1 = declare_device( - "_ZL17__floats2half2_rnff_1", - _type___half2(CPointer(float32), CPointer(float32)), + _ZL13__half2ull_rz6__half_nbst = declare_device( + "_ZL13__half2ull_rz6__half_nbst", uint64(CPointer(_type___half)) ) - def _ZL17__floats2half2_rnff_1_caller(arg_0, arg_1): - return _ZL17__floats2half2_rnff_1(arg_0, arg_1) + def _ZL13__half2ull_rz6__half_nbst_caller(arg_0): + return _ZL13__half2ull_rz6__half_nbst(arg_0) - @lower(__floats2half2_rn, float32, float32) + @lower(__half2ull_rz, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL17__floats2half2_rnff_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2ull_rz6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL17__floats2half2_rnff_1_caller, - signature(_type___half2, CPointer(float32), CPointer(float32)), + _ZL13__half2ull_rz6__half_nbst_caller, + signature(uint64, CPointer(_type___half)), ptrs, ) -_lower__ZL17__floats2half2_rnff_1(shim_stream, shim_obj) +_lower__ZL13__half2ull_rz6__half_nbst(shim_stream, shim_obj) -def __low2float(): +def __half2int_rn(): pass -def _lower__ZL11__low2float7__half2_1(shim_stream, shim_obj): +def _lower__ZL13__half2int_rn6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__low2float7__half2_1(float &retval , __half2* a) { - retval = __low2float(*a); + _ZL13__half2int_rn6__half_nbst(int &retval , __half* h) { + retval = __half2int_rn(*h); return 0; } """ - _ZL11__low2float7__half2_1 = declare_device( - "_ZL11__low2float7__half2_1", float32(CPointer(_type___half2)) + _ZL13__half2int_rn6__half_nbst = declare_device( + "_ZL13__half2int_rn6__half_nbst", int32(CPointer(_type___half)) ) - def _ZL11__low2float7__half2_1_caller(arg_0): - return _ZL11__low2float7__half2_1(arg_0) + def _ZL13__half2int_rn6__half_nbst_caller(arg_0): + return _ZL13__half2int_rn6__half_nbst(arg_0) - @lower(__low2float, _type___half2) + @lower(__half2int_rn, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__low2float7__half2_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2int_rn6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__low2float7__half2_1_caller, - signature(float32, CPointer(_type___half2)), + _ZL13__half2int_rn6__half_nbst_caller, + signature(int32, CPointer(_type___half)), ptrs, ) -_lower__ZL11__low2float7__half2_1(shim_stream, shim_obj) +_lower__ZL13__half2int_rn6__half_nbst(shim_stream, shim_obj) -def __high2float(): +def __half2int_rd(): pass -def _lower__ZL12__high2float7__half2_1(shim_stream, shim_obj): +def _lower__ZL13__half2int_rd6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__high2float7__half2_1(float &retval , __half2* a) { - retval = __high2float(*a); + _ZL13__half2int_rd6__half_nbst(int &retval , __half* h) { + retval = __half2int_rd(*h); return 0; } """ - _ZL12__high2float7__half2_1 = declare_device( - "_ZL12__high2float7__half2_1", float32(CPointer(_type___half2)) + _ZL13__half2int_rd6__half_nbst = declare_device( + "_ZL13__half2int_rd6__half_nbst", int32(CPointer(_type___half)) ) - def _ZL12__high2float7__half2_1_caller(arg_0): - return _ZL12__high2float7__half2_1(arg_0) + def _ZL13__half2int_rd6__half_nbst_caller(arg_0): + return _ZL13__half2int_rd6__half_nbst(arg_0) - @lower(__high2float, _type___half2) + @lower(__half2int_rd, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__high2float7__half2_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2int_rd6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__high2float7__half2_1_caller, - signature(float32, CPointer(_type___half2)), + _ZL13__half2int_rd6__half_nbst_caller, + signature(int32, CPointer(_type___half)), ptrs, ) -_lower__ZL12__high2float7__half2_1(shim_stream, shim_obj) +_lower__ZL13__half2int_rd6__half_nbst(shim_stream, shim_obj) -def __half2char_rz(): +def __half2int_ru(): pass -def _lower__ZL14__half2char_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL13__half2int_ru6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__half2char_rz6__half_1(signed char &retval , __half* h) { - retval = __half2char_rz(*h); + _ZL13__half2int_ru6__half_nbst(int &retval , __half* h) { + retval = __half2int_ru(*h); return 0; } """ - _ZL14__half2char_rz6__half_1 = declare_device( - "_ZL14__half2char_rz6__half_1", int8(CPointer(_type___half)) + _ZL13__half2int_ru6__half_nbst = declare_device( + "_ZL13__half2int_ru6__half_nbst", int32(CPointer(_type___half)) ) - def _ZL14__half2char_rz6__half_1_caller(arg_0): - return _ZL14__half2char_rz6__half_1(arg_0) + def _ZL13__half2int_ru6__half_nbst_caller(arg_0): + return _ZL13__half2int_ru6__half_nbst(arg_0) - @lower(__half2char_rz, _type___half) + @lower(__half2int_ru, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__half2char_rz6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2int_ru6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__half2char_rz6__half_1_caller, - signature(int8, CPointer(_type___half)), + _ZL13__half2int_ru6__half_nbst_caller, + signature(int32, CPointer(_type___half)), ptrs, ) -_lower__ZL14__half2char_rz6__half_1(shim_stream, shim_obj) +_lower__ZL13__half2int_ru6__half_nbst(shim_stream, shim_obj) -def __half2uchar_rz(): +def __int2half_rn(): pass -def _lower__ZL15__half2uchar_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL13__int2half_rni_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__half2uchar_rz6__half_1(unsigned char &retval , __half* h) { - retval = __half2uchar_rz(*h); + _ZL13__int2half_rni_nbst(__half &retval , int* i) { + retval = __int2half_rn(*i); return 0; } """ - _ZL15__half2uchar_rz6__half_1 = declare_device( - "_ZL15__half2uchar_rz6__half_1", uint8(CPointer(_type___half)) + _ZL13__int2half_rni_nbst = declare_device( + "_ZL13__int2half_rni_nbst", _type___half(CPointer(int32)) ) - def _ZL15__half2uchar_rz6__half_1_caller(arg_0): - return _ZL15__half2uchar_rz6__half_1(arg_0) + def _ZL13__int2half_rni_nbst_caller(arg_0): + return _ZL13__int2half_rni_nbst(arg_0) - @lower(__half2uchar_rz, _type___half) + @lower(__int2half_rn, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL15__half2uchar_rz6__half_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL13__int2half_rni_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__half2uchar_rz6__half_1_caller, - signature(uint8, CPointer(_type___half)), + _ZL13__int2half_rni_nbst_caller, + signature(_type___half, CPointer(int32)), ptrs, ) -_lower__ZL15__half2uchar_rz6__half_1(shim_stream, shim_obj) +_lower__ZL13__int2half_rni_nbst(shim_stream, shim_obj) -def __half2short_rz(): +def __int2half_rz(): pass -def _lower__ZL15__half2short_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL13__int2half_rzi_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__half2short_rz6__half_1(short &retval , __half* h) { - retval = __half2short_rz(*h); + _ZL13__int2half_rzi_nbst(__half &retval , int* i) { + retval = __int2half_rz(*i); return 0; } """ - _ZL15__half2short_rz6__half_1 = declare_device( - "_ZL15__half2short_rz6__half_1", int16(CPointer(_type___half)) + _ZL13__int2half_rzi_nbst = declare_device( + "_ZL13__int2half_rzi_nbst", _type___half(CPointer(int32)) ) - def _ZL15__half2short_rz6__half_1_caller(arg_0): - return _ZL15__half2short_rz6__half_1(arg_0) + def _ZL13__int2half_rzi_nbst_caller(arg_0): + return _ZL13__int2half_rzi_nbst(arg_0) - @lower(__half2short_rz, _type___half) + @lower(__int2half_rz, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL15__half2short_rz6__half_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL13__int2half_rzi_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__half2short_rz6__half_1_caller, - signature(int16, CPointer(_type___half)), + _ZL13__int2half_rzi_nbst_caller, + signature(_type___half, CPointer(int32)), ptrs, ) -_lower__ZL15__half2short_rz6__half_1(shim_stream, shim_obj) +_lower__ZL13__int2half_rzi_nbst(shim_stream, shim_obj) -def __half2ushort_rz(): +def __int2half_rd(): pass -def _lower__ZL16__half2ushort_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL13__int2half_rdi_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__half2ushort_rz6__half_1(unsigned short &retval , __half* h) { - retval = __half2ushort_rz(*h); + _ZL13__int2half_rdi_nbst(__half &retval , int* i) { + retval = __int2half_rd(*i); return 0; } """ - _ZL16__half2ushort_rz6__half_1 = declare_device( - "_ZL16__half2ushort_rz6__half_1", uint16(CPointer(_type___half)) + _ZL13__int2half_rdi_nbst = declare_device( + "_ZL13__int2half_rdi_nbst", _type___half(CPointer(int32)) ) - def _ZL16__half2ushort_rz6__half_1_caller(arg_0): - return _ZL16__half2ushort_rz6__half_1(arg_0) + def _ZL13__int2half_rdi_nbst_caller(arg_0): + return _ZL13__int2half_rdi_nbst(arg_0) - @lower(__half2ushort_rz, _type___half) + @lower(__int2half_rd, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL16__half2ushort_rz6__half_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL13__int2half_rdi_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__half2ushort_rz6__half_1_caller, - signature(uint16, CPointer(_type___half)), + _ZL13__int2half_rdi_nbst_caller, + signature(_type___half, CPointer(int32)), ptrs, ) -_lower__ZL16__half2ushort_rz6__half_1(shim_stream, shim_obj) +_lower__ZL13__int2half_rdi_nbst(shim_stream, shim_obj) -def __half2int_rz(): +def __int2half_ru(): pass -def _lower__ZL13__half2int_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL13__int2half_rui_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2int_rz6__half_1(int &retval , __half* h) { - retval = __half2int_rz(*h); + _ZL13__int2half_rui_nbst(__half &retval , int* i) { + retval = __int2half_ru(*i); return 0; } """ - _ZL13__half2int_rz6__half_1 = declare_device( - "_ZL13__half2int_rz6__half_1", int32(CPointer(_type___half)) + _ZL13__int2half_rui_nbst = declare_device( + "_ZL13__int2half_rui_nbst", _type___half(CPointer(int32)) ) - def _ZL13__half2int_rz6__half_1_caller(arg_0): - return _ZL13__half2int_rz6__half_1(arg_0) + def _ZL13__int2half_rui_nbst_caller(arg_0): + return _ZL13__int2half_rui_nbst(arg_0) - @lower(__half2int_rz, _type___half) + @lower(__int2half_ru, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2int_rz6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL13__int2half_rui_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2int_rz6__half_1_caller, - signature(int32, CPointer(_type___half)), + _ZL13__int2half_rui_nbst_caller, + signature(_type___half, CPointer(int32)), ptrs, ) -_lower__ZL13__half2int_rz6__half_1(shim_stream, shim_obj) +_lower__ZL13__int2half_rui_nbst(shim_stream, shim_obj) -def __half2uint_rz(): +def __half2short_rn(): pass -def _lower__ZL14__half2uint_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL15__half2short_rn6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__half2uint_rz6__half_1(unsigned int &retval , __half* h) { - retval = __half2uint_rz(*h); + _ZL15__half2short_rn6__half_nbst(short &retval , __half* h) { + retval = __half2short_rn(*h); return 0; } """ - _ZL14__half2uint_rz6__half_1 = declare_device( - "_ZL14__half2uint_rz6__half_1", uint32(CPointer(_type___half)) + _ZL15__half2short_rn6__half_nbst = declare_device( + "_ZL15__half2short_rn6__half_nbst", int16(CPointer(_type___half)) ) - def _ZL14__half2uint_rz6__half_1_caller(arg_0): - return _ZL14__half2uint_rz6__half_1(arg_0) + def _ZL15__half2short_rn6__half_nbst_caller(arg_0): + return _ZL15__half2short_rn6__half_nbst(arg_0) - @lower(__half2uint_rz, _type___half) + @lower(__half2short_rn, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__half2uint_rz6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL15__half2short_rn6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__half2uint_rz6__half_1_caller, - signature(uint32, CPointer(_type___half)), + _ZL15__half2short_rn6__half_nbst_caller, + signature(int16, CPointer(_type___half)), ptrs, ) -_lower__ZL14__half2uint_rz6__half_1(shim_stream, shim_obj) +_lower__ZL15__half2short_rn6__half_nbst(shim_stream, shim_obj) -def __half2ll_rz(): +def __half2short_rd(): pass -def _lower__ZL12__half2ll_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL15__half2short_rd6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__half2ll_rz6__half_1(long long &retval , __half* h) { - retval = __half2ll_rz(*h); + _ZL15__half2short_rd6__half_nbst(short &retval , __half* h) { + retval = __half2short_rd(*h); return 0; } """ - _ZL12__half2ll_rz6__half_1 = declare_device( - "_ZL12__half2ll_rz6__half_1", int64(CPointer(_type___half)) + _ZL15__half2short_rd6__half_nbst = declare_device( + "_ZL15__half2short_rd6__half_nbst", int16(CPointer(_type___half)) ) - def _ZL12__half2ll_rz6__half_1_caller(arg_0): - return _ZL12__half2ll_rz6__half_1(arg_0) + def _ZL15__half2short_rd6__half_nbst_caller(arg_0): + return _ZL15__half2short_rd6__half_nbst(arg_0) - @lower(__half2ll_rz, _type___half) + @lower(__half2short_rd, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__half2ll_rz6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL15__half2short_rd6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__half2ll_rz6__half_1_caller, - signature(int64, CPointer(_type___half)), + _ZL15__half2short_rd6__half_nbst_caller, + signature(int16, CPointer(_type___half)), ptrs, ) -_lower__ZL12__half2ll_rz6__half_1(shim_stream, shim_obj) +_lower__ZL15__half2short_rd6__half_nbst(shim_stream, shim_obj) -def __half2ull_rz(): +def __half2short_ru(): pass -def _lower__ZL13__half2ull_rz6__half_1(shim_stream, shim_obj): +def _lower__ZL15__half2short_ru6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2ull_rz6__half_1(unsigned long long &retval , __half* h) { - retval = __half2ull_rz(*h); + _ZL15__half2short_ru6__half_nbst(short &retval , __half* h) { + retval = __half2short_ru(*h); return 0; } """ - _ZL13__half2ull_rz6__half_1 = declare_device( - "_ZL13__half2ull_rz6__half_1", uint64(CPointer(_type___half)) + _ZL15__half2short_ru6__half_nbst = declare_device( + "_ZL15__half2short_ru6__half_nbst", int16(CPointer(_type___half)) ) - def _ZL13__half2ull_rz6__half_1_caller(arg_0): - return _ZL13__half2ull_rz6__half_1(arg_0) + def _ZL15__half2short_ru6__half_nbst_caller(arg_0): + return _ZL15__half2short_ru6__half_nbst(arg_0) - @lower(__half2ull_rz, _type___half) + @lower(__half2short_ru, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2ull_rz6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL15__half2short_ru6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2ull_rz6__half_1_caller, - signature(uint64, CPointer(_type___half)), + _ZL15__half2short_ru6__half_nbst_caller, + signature(int16, CPointer(_type___half)), ptrs, ) -_lower__ZL13__half2ull_rz6__half_1(shim_stream, shim_obj) +_lower__ZL15__half2short_ru6__half_nbst(shim_stream, shim_obj) -def make_half2(): +def __short2half_rn(): pass -def _lower__ZL10make_half26__halfS__1(shim_stream, shim_obj): +def _lower__ZL15__short2half_rns_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL10make_half26__halfS__1(__half2 &retval , __half* x, __half* y) { - retval = make_half2(*x, *y); + _ZL15__short2half_rns_nbst(__half &retval , short* i) { + retval = __short2half_rn(*i); return 0; } """ - _ZL10make_half26__halfS__1 = declare_device( - "_ZL10make_half26__halfS__1", - _type___half2(CPointer(_type___half), CPointer(_type___half)), + _ZL15__short2half_rns_nbst = declare_device( + "_ZL15__short2half_rns_nbst", _type___half(CPointer(int16)) ) - def _ZL10make_half26__halfS__1_caller(arg_0, arg_1): - return _ZL10make_half26__halfS__1(arg_0, arg_1) + def _ZL15__short2half_rns_nbst_caller(arg_0): + return _ZL15__short2half_rns_nbst(arg_0) - @lower(make_half2, _type___half, _type___half) + @lower(__short2half_rn, int16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10make_half26__halfS__1", shim_raw_str) + shim_stream.write_with_key("_ZL15__short2half_rns_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL10make_half26__halfS__1_caller, - signature( - _type___half2, CPointer(_type___half), CPointer(_type___half) - ), + _ZL15__short2half_rns_nbst_caller, + signature(_type___half, CPointer(int16)), ptrs, ) -_lower__ZL10make_half26__halfS__1(shim_stream, shim_obj) +_lower__ZL15__short2half_rns_nbst(shim_stream, shim_obj) -def __float22half2_rn(): +def __short2half_rz(): pass -def _lower__ZL17__float22half2_rn6float2_1(shim_stream, shim_obj): +def _lower__ZL15__short2half_rzs_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL17__float22half2_rn6float2_1(__half2 &retval , float2* a) { - retval = __float22half2_rn(*a); + _ZL15__short2half_rzs_nbst(__half &retval , short* i) { + retval = __short2half_rz(*i); return 0; } """ - _ZL17__float22half2_rn6float2_1 = declare_device( - "_ZL17__float22half2_rn6float2_1", _type___half2(CPointer(float32x2)) + _ZL15__short2half_rzs_nbst = declare_device( + "_ZL15__short2half_rzs_nbst", _type___half(CPointer(int16)) ) - def _ZL17__float22half2_rn6float2_1_caller(arg_0): - return _ZL17__float22half2_rn6float2_1(arg_0) + def _ZL15__short2half_rzs_nbst_caller(arg_0): + return _ZL15__short2half_rzs_nbst(arg_0) - @lower(__float22half2_rn, float32x2) + @lower(__short2half_rz, int16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL17__float22half2_rn6float2_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL15__short2half_rzs_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL17__float22half2_rn6float2_1_caller, - signature(_type___half2, CPointer(float32x2)), + _ZL15__short2half_rzs_nbst_caller, + signature(_type___half, CPointer(int16)), ptrs, ) -_lower__ZL17__float22half2_rn6float2_1(shim_stream, shim_obj) +_lower__ZL15__short2half_rzs_nbst(shim_stream, shim_obj) -def __half22float2(): +def __short2half_rd(): pass -def _lower__ZL14__half22float27__half2_1(shim_stream, shim_obj): +def _lower__ZL15__short2half_rds_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__half22float27__half2_1(float2 &retval , __half2* a) { - retval = __half22float2(*a); + _ZL15__short2half_rds_nbst(__half &retval , short* i) { + retval = __short2half_rd(*i); return 0; } """ - _ZL14__half22float27__half2_1 = declare_device( - "_ZL14__half22float27__half2_1", float32x2(CPointer(_type___half2)) + _ZL15__short2half_rds_nbst = declare_device( + "_ZL15__short2half_rds_nbst", _type___half(CPointer(int16)) ) - def _ZL14__half22float27__half2_1_caller(arg_0): - return _ZL14__half22float27__half2_1(arg_0) + def _ZL15__short2half_rds_nbst_caller(arg_0): + return _ZL15__short2half_rds_nbst(arg_0) - @lower(__half22float2, _type___half2) + @lower(__short2half_rd, int16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL14__half22float27__half2_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL15__short2half_rds_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__half22float27__half2_1_caller, - signature(float32x2, CPointer(_type___half2)), + _ZL15__short2half_rds_nbst_caller, + signature(_type___half, CPointer(int16)), ptrs, ) -_lower__ZL14__half22float27__half2_1(shim_stream, shim_obj) +_lower__ZL15__short2half_rds_nbst(shim_stream, shim_obj) -def __half2int_rn(): +def __short2half_ru(): pass -def _lower__ZL13__half2int_rn6__half_1(shim_stream, shim_obj): +def _lower__ZL15__short2half_rus_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2int_rn6__half_1(int &retval , __half* h) { - retval = __half2int_rn(*h); + _ZL15__short2half_rus_nbst(__half &retval , short* i) { + retval = __short2half_ru(*i); return 0; } """ - _ZL13__half2int_rn6__half_1 = declare_device( - "_ZL13__half2int_rn6__half_1", int32(CPointer(_type___half)) + _ZL15__short2half_rus_nbst = declare_device( + "_ZL15__short2half_rus_nbst", _type___half(CPointer(int16)) ) - def _ZL13__half2int_rn6__half_1_caller(arg_0): - return _ZL13__half2int_rn6__half_1(arg_0) + def _ZL15__short2half_rus_nbst_caller(arg_0): + return _ZL15__short2half_rus_nbst(arg_0) - @lower(__half2int_rn, _type___half) + @lower(__short2half_ru, int16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2int_rn6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL15__short2half_rus_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2int_rn6__half_1_caller, - signature(int32, CPointer(_type___half)), + _ZL15__short2half_rus_nbst_caller, + signature(_type___half, CPointer(int16)), ptrs, ) -_lower__ZL13__half2int_rn6__half_1(shim_stream, shim_obj) +_lower__ZL15__short2half_rus_nbst(shim_stream, shim_obj) -def __half2int_rd(): +def __half2uint_rn(): pass -def _lower__ZL13__half2int_rd6__half_1(shim_stream, shim_obj): +def _lower__ZL14__half2uint_rn6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2int_rd6__half_1(int &retval , __half* h) { - retval = __half2int_rd(*h); + _ZL14__half2uint_rn6__half_nbst(unsigned int &retval , __half* h) { + retval = __half2uint_rn(*h); return 0; } """ - _ZL13__half2int_rd6__half_1 = declare_device( - "_ZL13__half2int_rd6__half_1", int32(CPointer(_type___half)) + _ZL14__half2uint_rn6__half_nbst = declare_device( + "_ZL14__half2uint_rn6__half_nbst", uint32(CPointer(_type___half)) ) - def _ZL13__half2int_rd6__half_1_caller(arg_0): - return _ZL13__half2int_rd6__half_1(arg_0) + def _ZL14__half2uint_rn6__half_nbst_caller(arg_0): + return _ZL14__half2uint_rn6__half_nbst(arg_0) - @lower(__half2int_rd, _type___half) + @lower(__half2uint_rn, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2int_rd6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL14__half2uint_rn6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2int_rd6__half_1_caller, - signature(int32, CPointer(_type___half)), + _ZL14__half2uint_rn6__half_nbst_caller, + signature(uint32, CPointer(_type___half)), ptrs, ) -_lower__ZL13__half2int_rd6__half_1(shim_stream, shim_obj) +_lower__ZL14__half2uint_rn6__half_nbst(shim_stream, shim_obj) -def __half2int_ru(): +def __half2uint_rd(): pass -def _lower__ZL13__half2int_ru6__half_1(shim_stream, shim_obj): +def _lower__ZL14__half2uint_rd6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2int_ru6__half_1(int &retval , __half* h) { - retval = __half2int_ru(*h); + _ZL14__half2uint_rd6__half_nbst(unsigned int &retval , __half* h) { + retval = __half2uint_rd(*h); return 0; } """ - _ZL13__half2int_ru6__half_1 = declare_device( - "_ZL13__half2int_ru6__half_1", int32(CPointer(_type___half)) + _ZL14__half2uint_rd6__half_nbst = declare_device( + "_ZL14__half2uint_rd6__half_nbst", uint32(CPointer(_type___half)) ) - def _ZL13__half2int_ru6__half_1_caller(arg_0): - return _ZL13__half2int_ru6__half_1(arg_0) + def _ZL14__half2uint_rd6__half_nbst_caller(arg_0): + return _ZL14__half2uint_rd6__half_nbst(arg_0) - @lower(__half2int_ru, _type___half) + @lower(__half2uint_rd, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2int_ru6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL14__half2uint_rd6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2int_ru6__half_1_caller, - signature(int32, CPointer(_type___half)), + _ZL14__half2uint_rd6__half_nbst_caller, + signature(uint32, CPointer(_type___half)), ptrs, ) -_lower__ZL13__half2int_ru6__half_1(shim_stream, shim_obj) +_lower__ZL14__half2uint_rd6__half_nbst(shim_stream, shim_obj) -def __int2half_rn(): +def __half2uint_ru(): pass -def _lower__ZL13__int2half_rni_1(shim_stream, shim_obj): +def _lower__ZL14__half2uint_ru6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__int2half_rni_1(__half &retval , int* i) { - retval = __int2half_rn(*i); + _ZL14__half2uint_ru6__half_nbst(unsigned int &retval , __half* h) { + retval = __half2uint_ru(*h); return 0; } """ - _ZL13__int2half_rni_1 = declare_device( - "_ZL13__int2half_rni_1", _type___half(CPointer(int32)) + _ZL14__half2uint_ru6__half_nbst = declare_device( + "_ZL14__half2uint_ru6__half_nbst", uint32(CPointer(_type___half)) ) - def _ZL13__int2half_rni_1_caller(arg_0): - return _ZL13__int2half_rni_1(arg_0) + def _ZL14__half2uint_ru6__half_nbst_caller(arg_0): + return _ZL14__half2uint_ru6__half_nbst(arg_0) - @lower(__int2half_rn, int32) + @lower(__half2uint_ru, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__int2half_rni_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL14__half2uint_ru6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__int2half_rni_1_caller, - signature(_type___half, CPointer(int32)), + _ZL14__half2uint_ru6__half_nbst_caller, + signature(uint32, CPointer(_type___half)), ptrs, ) -_lower__ZL13__int2half_rni_1(shim_stream, shim_obj) +_lower__ZL14__half2uint_ru6__half_nbst(shim_stream, shim_obj) -def __int2half_rz(): +def __uint2half_rn(): pass -def _lower__ZL13__int2half_rzi_1(shim_stream, shim_obj): +def _lower__ZL14__uint2half_rnj_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__int2half_rzi_1(__half &retval , int* i) { - retval = __int2half_rz(*i); + _ZL14__uint2half_rnj_nbst(__half &retval , unsigned int* i) { + retval = __uint2half_rn(*i); return 0; } """ - _ZL13__int2half_rzi_1 = declare_device( - "_ZL13__int2half_rzi_1", _type___half(CPointer(int32)) + _ZL14__uint2half_rnj_nbst = declare_device( + "_ZL14__uint2half_rnj_nbst", _type___half(CPointer(uint32)) ) - def _ZL13__int2half_rzi_1_caller(arg_0): - return _ZL13__int2half_rzi_1(arg_0) + def _ZL14__uint2half_rnj_nbst_caller(arg_0): + return _ZL14__uint2half_rnj_nbst(arg_0) - @lower(__int2half_rz, int32) + @lower(__uint2half_rn, uint32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__int2half_rzi_1", shim_raw_str) + shim_stream.write_with_key("_ZL14__uint2half_rnj_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__int2half_rzi_1_caller, - signature(_type___half, CPointer(int32)), + _ZL14__uint2half_rnj_nbst_caller, + signature(_type___half, CPointer(uint32)), ptrs, ) -_lower__ZL13__int2half_rzi_1(shim_stream, shim_obj) +_lower__ZL14__uint2half_rnj_nbst(shim_stream, shim_obj) -def __int2half_rd(): +def __uint2half_rz(): pass -def _lower__ZL13__int2half_rdi_1(shim_stream, shim_obj): +def _lower__ZL14__uint2half_rzj_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__int2half_rdi_1(__half &retval , int* i) { - retval = __int2half_rd(*i); + _ZL14__uint2half_rzj_nbst(__half &retval , unsigned int* i) { + retval = __uint2half_rz(*i); return 0; } """ - _ZL13__int2half_rdi_1 = declare_device( - "_ZL13__int2half_rdi_1", _type___half(CPointer(int32)) + _ZL14__uint2half_rzj_nbst = declare_device( + "_ZL14__uint2half_rzj_nbst", _type___half(CPointer(uint32)) ) - def _ZL13__int2half_rdi_1_caller(arg_0): - return _ZL13__int2half_rdi_1(arg_0) + def _ZL14__uint2half_rzj_nbst_caller(arg_0): + return _ZL14__uint2half_rzj_nbst(arg_0) - @lower(__int2half_rd, int32) + @lower(__uint2half_rz, uint32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__int2half_rdi_1", shim_raw_str) + shim_stream.write_with_key("_ZL14__uint2half_rzj_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__int2half_rdi_1_caller, - signature(_type___half, CPointer(int32)), + _ZL14__uint2half_rzj_nbst_caller, + signature(_type___half, CPointer(uint32)), ptrs, ) -_lower__ZL13__int2half_rdi_1(shim_stream, shim_obj) +_lower__ZL14__uint2half_rzj_nbst(shim_stream, shim_obj) -def __int2half_ru(): +def __uint2half_rd(): pass -def _lower__ZL13__int2half_rui_1(shim_stream, shim_obj): +def _lower__ZL14__uint2half_rdj_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__int2half_rui_1(__half &retval , int* i) { - retval = __int2half_ru(*i); + _ZL14__uint2half_rdj_nbst(__half &retval , unsigned int* i) { + retval = __uint2half_rd(*i); return 0; } """ - _ZL13__int2half_rui_1 = declare_device( - "_ZL13__int2half_rui_1", _type___half(CPointer(int32)) + _ZL14__uint2half_rdj_nbst = declare_device( + "_ZL14__uint2half_rdj_nbst", _type___half(CPointer(uint32)) ) - def _ZL13__int2half_rui_1_caller(arg_0): - return _ZL13__int2half_rui_1(arg_0) + def _ZL14__uint2half_rdj_nbst_caller(arg_0): + return _ZL14__uint2half_rdj_nbst(arg_0) - @lower(__int2half_ru, int32) + @lower(__uint2half_rd, uint32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__int2half_rui_1", shim_raw_str) + shim_stream.write_with_key("_ZL14__uint2half_rdj_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__int2half_rui_1_caller, - signature(_type___half, CPointer(int32)), + _ZL14__uint2half_rdj_nbst_caller, + signature(_type___half, CPointer(uint32)), ptrs, ) -_lower__ZL13__int2half_rui_1(shim_stream, shim_obj) +_lower__ZL14__uint2half_rdj_nbst(shim_stream, shim_obj) -def __half2short_rn(): +def __uint2half_ru(): pass -def _lower__ZL15__half2short_rn6__half_1(shim_stream, shim_obj): +def _lower__ZL14__uint2half_ruj_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__half2short_rn6__half_1(short &retval , __half* h) { - retval = __half2short_rn(*h); + _ZL14__uint2half_ruj_nbst(__half &retval , unsigned int* i) { + retval = __uint2half_ru(*i); return 0; } """ - _ZL15__half2short_rn6__half_1 = declare_device( - "_ZL15__half2short_rn6__half_1", int16(CPointer(_type___half)) + _ZL14__uint2half_ruj_nbst = declare_device( + "_ZL14__uint2half_ruj_nbst", _type___half(CPointer(uint32)) ) - def _ZL15__half2short_rn6__half_1_caller(arg_0): - return _ZL15__half2short_rn6__half_1(arg_0) + def _ZL14__uint2half_ruj_nbst_caller(arg_0): + return _ZL14__uint2half_ruj_nbst(arg_0) - @lower(__half2short_rn, _type___half) + @lower(__uint2half_ru, uint32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL15__half2short_rn6__half_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL14__uint2half_ruj_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__half2short_rn6__half_1_caller, - signature(int16, CPointer(_type___half)), + _ZL14__uint2half_ruj_nbst_caller, + signature(_type___half, CPointer(uint32)), ptrs, ) -_lower__ZL15__half2short_rn6__half_1(shim_stream, shim_obj) +_lower__ZL14__uint2half_ruj_nbst(shim_stream, shim_obj) -def __half2short_rd(): +def __half2ushort_rn(): pass -def _lower__ZL15__half2short_rd6__half_1(shim_stream, shim_obj): +def _lower__ZL16__half2ushort_rn6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__half2short_rd6__half_1(short &retval , __half* h) { - retval = __half2short_rd(*h); + _ZL16__half2ushort_rn6__half_nbst(unsigned short &retval , __half* h) { + retval = __half2ushort_rn(*h); return 0; } """ - _ZL15__half2short_rd6__half_1 = declare_device( - "_ZL15__half2short_rd6__half_1", int16(CPointer(_type___half)) + _ZL16__half2ushort_rn6__half_nbst = declare_device( + "_ZL16__half2ushort_rn6__half_nbst", uint16(CPointer(_type___half)) ) - def _ZL15__half2short_rd6__half_1_caller(arg_0): - return _ZL15__half2short_rd6__half_1(arg_0) + def _ZL16__half2ushort_rn6__half_nbst_caller(arg_0): + return _ZL16__half2ushort_rn6__half_nbst(arg_0) - @lower(__half2short_rd, _type___half) + @lower(__half2ushort_rn, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) shim_stream.write_with_key( - "_ZL15__half2short_rd6__half_1", shim_raw_str + "_ZL16__half2ushort_rn6__half_nbst", shim_raw_str ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): @@ -2433,40 +2391,40 @@ def impl(context, builder, sig, args): return context.compile_internal( builder, - _ZL15__half2short_rd6__half_1_caller, - signature(int16, CPointer(_type___half)), + _ZL16__half2ushort_rn6__half_nbst_caller, + signature(uint16, CPointer(_type___half)), ptrs, ) -_lower__ZL15__half2short_rd6__half_1(shim_stream, shim_obj) +_lower__ZL16__half2ushort_rn6__half_nbst(shim_stream, shim_obj) -def __half2short_ru(): +def __half2ushort_rd(): pass -def _lower__ZL15__half2short_ru6__half_1(shim_stream, shim_obj): +def _lower__ZL16__half2ushort_rd6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__half2short_ru6__half_1(short &retval , __half* h) { - retval = __half2short_ru(*h); + _ZL16__half2ushort_rd6__half_nbst(unsigned short &retval , __half* h) { + retval = __half2ushort_rd(*h); return 0; } """ - _ZL15__half2short_ru6__half_1 = declare_device( - "_ZL15__half2short_ru6__half_1", int16(CPointer(_type___half)) + _ZL16__half2ushort_rd6__half_nbst = declare_device( + "_ZL16__half2ushort_rd6__half_nbst", uint16(CPointer(_type___half)) ) - def _ZL15__half2short_ru6__half_1_caller(arg_0): - return _ZL15__half2short_ru6__half_1(arg_0) + def _ZL16__half2ushort_rd6__half_nbst_caller(arg_0): + return _ZL16__half2ushort_rd6__half_nbst(arg_0) - @lower(__half2short_ru, _type___half) + @lower(__half2ushort_rd, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) shim_stream.write_with_key( - "_ZL15__half2short_ru6__half_1", shim_raw_str + "_ZL16__half2ushort_rd6__half_nbst", shim_raw_str ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): @@ -2474,510 +2432,516 @@ def impl(context, builder, sig, args): return context.compile_internal( builder, - _ZL15__half2short_ru6__half_1_caller, - signature(int16, CPointer(_type___half)), + _ZL16__half2ushort_rd6__half_nbst_caller, + signature(uint16, CPointer(_type___half)), ptrs, ) -_lower__ZL15__half2short_ru6__half_1(shim_stream, shim_obj) +_lower__ZL16__half2ushort_rd6__half_nbst(shim_stream, shim_obj) -def __short2half_rn(): +def __half2ushort_ru(): pass -def _lower__ZL15__short2half_rns_1(shim_stream, shim_obj): +def _lower__ZL16__half2ushort_ru6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__short2half_rns_1(__half &retval , short* i) { - retval = __short2half_rn(*i); + _ZL16__half2ushort_ru6__half_nbst(unsigned short &retval , __half* h) { + retval = __half2ushort_ru(*h); return 0; } """ - _ZL15__short2half_rns_1 = declare_device( - "_ZL15__short2half_rns_1", _type___half(CPointer(int16)) + _ZL16__half2ushort_ru6__half_nbst = declare_device( + "_ZL16__half2ushort_ru6__half_nbst", uint16(CPointer(_type___half)) ) - def _ZL15__short2half_rns_1_caller(arg_0): - return _ZL15__short2half_rns_1(arg_0) + def _ZL16__half2ushort_ru6__half_nbst_caller(arg_0): + return _ZL16__half2ushort_ru6__half_nbst(arg_0) - @lower(__short2half_rn, int16) + @lower(__half2ushort_ru, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__short2half_rns_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL16__half2ushort_ru6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__short2half_rns_1_caller, - signature(_type___half, CPointer(int16)), + _ZL16__half2ushort_ru6__half_nbst_caller, + signature(uint16, CPointer(_type___half)), ptrs, ) -_lower__ZL15__short2half_rns_1(shim_stream, shim_obj) +_lower__ZL16__half2ushort_ru6__half_nbst(shim_stream, shim_obj) -def __short2half_rz(): +def __ushort2half_rn(): pass -def _lower__ZL15__short2half_rzs_1(shim_stream, shim_obj): +def _lower__ZL16__ushort2half_rnt_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__short2half_rzs_1(__half &retval , short* i) { - retval = __short2half_rz(*i); + _ZL16__ushort2half_rnt_nbst(__half &retval , unsigned short* i) { + retval = __ushort2half_rn(*i); return 0; } """ - _ZL15__short2half_rzs_1 = declare_device( - "_ZL15__short2half_rzs_1", _type___half(CPointer(int16)) + _ZL16__ushort2half_rnt_nbst = declare_device( + "_ZL16__ushort2half_rnt_nbst", _type___half(CPointer(uint16)) ) - def _ZL15__short2half_rzs_1_caller(arg_0): - return _ZL15__short2half_rzs_1(arg_0) + def _ZL16__ushort2half_rnt_nbst_caller(arg_0): + return _ZL16__ushort2half_rnt_nbst(arg_0) - @lower(__short2half_rz, int16) + @lower(__ushort2half_rn, uint16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__short2half_rzs_1", shim_raw_str) + shim_stream.write_with_key("_ZL16__ushort2half_rnt_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__short2half_rzs_1_caller, - signature(_type___half, CPointer(int16)), + _ZL16__ushort2half_rnt_nbst_caller, + signature(_type___half, CPointer(uint16)), ptrs, ) -_lower__ZL15__short2half_rzs_1(shim_stream, shim_obj) +_lower__ZL16__ushort2half_rnt_nbst(shim_stream, shim_obj) -def __short2half_rd(): +def __ushort2half_rz(): pass -def _lower__ZL15__short2half_rds_1(shim_stream, shim_obj): +def _lower__ZL16__ushort2half_rzt_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__short2half_rds_1(__half &retval , short* i) { - retval = __short2half_rd(*i); + _ZL16__ushort2half_rzt_nbst(__half &retval , unsigned short* i) { + retval = __ushort2half_rz(*i); return 0; } """ - _ZL15__short2half_rds_1 = declare_device( - "_ZL15__short2half_rds_1", _type___half(CPointer(int16)) + _ZL16__ushort2half_rzt_nbst = declare_device( + "_ZL16__ushort2half_rzt_nbst", _type___half(CPointer(uint16)) ) - def _ZL15__short2half_rds_1_caller(arg_0): - return _ZL15__short2half_rds_1(arg_0) + def _ZL16__ushort2half_rzt_nbst_caller(arg_0): + return _ZL16__ushort2half_rzt_nbst(arg_0) - @lower(__short2half_rd, int16) + @lower(__ushort2half_rz, uint16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__short2half_rds_1", shim_raw_str) + shim_stream.write_with_key("_ZL16__ushort2half_rzt_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__short2half_rds_1_caller, - signature(_type___half, CPointer(int16)), + _ZL16__ushort2half_rzt_nbst_caller, + signature(_type___half, CPointer(uint16)), ptrs, ) -_lower__ZL15__short2half_rds_1(shim_stream, shim_obj) +_lower__ZL16__ushort2half_rzt_nbst(shim_stream, shim_obj) -def __short2half_ru(): +def __ushort2half_rd(): pass -def _lower__ZL15__short2half_rus_1(shim_stream, shim_obj): +def _lower__ZL16__ushort2half_rdt_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__short2half_rus_1(__half &retval , short* i) { - retval = __short2half_ru(*i); + _ZL16__ushort2half_rdt_nbst(__half &retval , unsigned short* i) { + retval = __ushort2half_rd(*i); return 0; } """ - _ZL15__short2half_rus_1 = declare_device( - "_ZL15__short2half_rus_1", _type___half(CPointer(int16)) + _ZL16__ushort2half_rdt_nbst = declare_device( + "_ZL16__ushort2half_rdt_nbst", _type___half(CPointer(uint16)) ) - def _ZL15__short2half_rus_1_caller(arg_0): - return _ZL15__short2half_rus_1(arg_0) + def _ZL16__ushort2half_rdt_nbst_caller(arg_0): + return _ZL16__ushort2half_rdt_nbst(arg_0) - @lower(__short2half_ru, int16) + @lower(__ushort2half_rd, uint16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__short2half_rus_1", shim_raw_str) + shim_stream.write_with_key("_ZL16__ushort2half_rdt_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__short2half_rus_1_caller, - signature(_type___half, CPointer(int16)), + _ZL16__ushort2half_rdt_nbst_caller, + signature(_type___half, CPointer(uint16)), ptrs, ) -_lower__ZL15__short2half_rus_1(shim_stream, shim_obj) +_lower__ZL16__ushort2half_rdt_nbst(shim_stream, shim_obj) -def __half2uint_rn(): +def __ushort2half_ru(): pass -def _lower__ZL14__half2uint_rn6__half_1(shim_stream, shim_obj): +def _lower__ZL16__ushort2half_rut_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__half2uint_rn6__half_1(unsigned int &retval , __half* h) { - retval = __half2uint_rn(*h); + _ZL16__ushort2half_rut_nbst(__half &retval , unsigned short* i) { + retval = __ushort2half_ru(*i); return 0; } """ - _ZL14__half2uint_rn6__half_1 = declare_device( - "_ZL14__half2uint_rn6__half_1", uint32(CPointer(_type___half)) + _ZL16__ushort2half_rut_nbst = declare_device( + "_ZL16__ushort2half_rut_nbst", _type___half(CPointer(uint16)) ) - def _ZL14__half2uint_rn6__half_1_caller(arg_0): - return _ZL14__half2uint_rn6__half_1(arg_0) + def _ZL16__ushort2half_rut_nbst_caller(arg_0): + return _ZL16__ushort2half_rut_nbst(arg_0) - @lower(__half2uint_rn, _type___half) + @lower(__ushort2half_ru, uint16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__half2uint_rn6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL16__ushort2half_rut_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__half2uint_rn6__half_1_caller, - signature(uint32, CPointer(_type___half)), + _ZL16__ushort2half_rut_nbst_caller, + signature(_type___half, CPointer(uint16)), ptrs, ) -_lower__ZL14__half2uint_rn6__half_1(shim_stream, shim_obj) +_lower__ZL16__ushort2half_rut_nbst(shim_stream, shim_obj) -def __half2uint_rd(): +def __half2ull_rn(): pass -def _lower__ZL14__half2uint_rd6__half_1(shim_stream, shim_obj): +def _lower__ZL13__half2ull_rn6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__half2uint_rd6__half_1(unsigned int &retval , __half* h) { - retval = __half2uint_rd(*h); + _ZL13__half2ull_rn6__half_nbst(unsigned long long &retval , __half* h) { + retval = __half2ull_rn(*h); return 0; } """ - _ZL14__half2uint_rd6__half_1 = declare_device( - "_ZL14__half2uint_rd6__half_1", uint32(CPointer(_type___half)) + _ZL13__half2ull_rn6__half_nbst = declare_device( + "_ZL13__half2ull_rn6__half_nbst", uint64(CPointer(_type___half)) ) - def _ZL14__half2uint_rd6__half_1_caller(arg_0): - return _ZL14__half2uint_rd6__half_1(arg_0) + def _ZL13__half2ull_rn6__half_nbst_caller(arg_0): + return _ZL13__half2ull_rn6__half_nbst(arg_0) - @lower(__half2uint_rd, _type___half) + @lower(__half2ull_rn, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__half2uint_rd6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2ull_rn6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__half2uint_rd6__half_1_caller, - signature(uint32, CPointer(_type___half)), + _ZL13__half2ull_rn6__half_nbst_caller, + signature(uint64, CPointer(_type___half)), ptrs, ) -_lower__ZL14__half2uint_rd6__half_1(shim_stream, shim_obj) +_lower__ZL13__half2ull_rn6__half_nbst(shim_stream, shim_obj) -def __half2uint_ru(): +def __half2ull_rd(): pass -def _lower__ZL14__half2uint_ru6__half_1(shim_stream, shim_obj): +def _lower__ZL13__half2ull_rd6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__half2uint_ru6__half_1(unsigned int &retval , __half* h) { - retval = __half2uint_ru(*h); + _ZL13__half2ull_rd6__half_nbst(unsigned long long &retval , __half* h) { + retval = __half2ull_rd(*h); return 0; } """ - _ZL14__half2uint_ru6__half_1 = declare_device( - "_ZL14__half2uint_ru6__half_1", uint32(CPointer(_type___half)) + _ZL13__half2ull_rd6__half_nbst = declare_device( + "_ZL13__half2ull_rd6__half_nbst", uint64(CPointer(_type___half)) ) - def _ZL14__half2uint_ru6__half_1_caller(arg_0): - return _ZL14__half2uint_ru6__half_1(arg_0) + def _ZL13__half2ull_rd6__half_nbst_caller(arg_0): + return _ZL13__half2ull_rd6__half_nbst(arg_0) - @lower(__half2uint_ru, _type___half) + @lower(__half2ull_rd, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__half2uint_ru6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2ull_rd6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__half2uint_ru6__half_1_caller, - signature(uint32, CPointer(_type___half)), + _ZL13__half2ull_rd6__half_nbst_caller, + signature(uint64, CPointer(_type___half)), ptrs, ) -_lower__ZL14__half2uint_ru6__half_1(shim_stream, shim_obj) +_lower__ZL13__half2ull_rd6__half_nbst(shim_stream, shim_obj) -def __uint2half_rn(): +def __half2ull_ru(): pass -def _lower__ZL14__uint2half_rnj_1(shim_stream, shim_obj): +def _lower__ZL13__half2ull_ru6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__uint2half_rnj_1(__half &retval , unsigned int* i) { - retval = __uint2half_rn(*i); + _ZL13__half2ull_ru6__half_nbst(unsigned long long &retval , __half* h) { + retval = __half2ull_ru(*h); return 0; } """ - _ZL14__uint2half_rnj_1 = declare_device( - "_ZL14__uint2half_rnj_1", _type___half(CPointer(uint32)) + _ZL13__half2ull_ru6__half_nbst = declare_device( + "_ZL13__half2ull_ru6__half_nbst", uint64(CPointer(_type___half)) ) - def _ZL14__uint2half_rnj_1_caller(arg_0): - return _ZL14__uint2half_rnj_1(arg_0) + def _ZL13__half2ull_ru6__half_nbst_caller(arg_0): + return _ZL13__half2ull_ru6__half_nbst(arg_0) - @lower(__uint2half_rn, uint32) + @lower(__half2ull_ru, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__uint2half_rnj_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL13__half2ull_ru6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__uint2half_rnj_1_caller, - signature(_type___half, CPointer(uint32)), + _ZL13__half2ull_ru6__half_nbst_caller, + signature(uint64, CPointer(_type___half)), ptrs, ) -_lower__ZL14__uint2half_rnj_1(shim_stream, shim_obj) +_lower__ZL13__half2ull_ru6__half_nbst(shim_stream, shim_obj) -def __uint2half_rz(): +def __ull2half_rn(): pass -def _lower__ZL14__uint2half_rzj_1(shim_stream, shim_obj): +def _lower__ZL13__ull2half_rny_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__uint2half_rzj_1(__half &retval , unsigned int* i) { - retval = __uint2half_rz(*i); + _ZL13__ull2half_rny_nbst(__half &retval , unsigned long long* i) { + retval = __ull2half_rn(*i); return 0; } """ - _ZL14__uint2half_rzj_1 = declare_device( - "_ZL14__uint2half_rzj_1", _type___half(CPointer(uint32)) + _ZL13__ull2half_rny_nbst = declare_device( + "_ZL13__ull2half_rny_nbst", _type___half(CPointer(uint64)) ) - def _ZL14__uint2half_rzj_1_caller(arg_0): - return _ZL14__uint2half_rzj_1(arg_0) + def _ZL13__ull2half_rny_nbst_caller(arg_0): + return _ZL13__ull2half_rny_nbst(arg_0) - @lower(__uint2half_rz, uint32) + @lower(__ull2half_rn, uint64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__uint2half_rzj_1", shim_raw_str) + shim_stream.write_with_key("_ZL13__ull2half_rny_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__uint2half_rzj_1_caller, - signature(_type___half, CPointer(uint32)), + _ZL13__ull2half_rny_nbst_caller, + signature(_type___half, CPointer(uint64)), ptrs, ) -_lower__ZL14__uint2half_rzj_1(shim_stream, shim_obj) +_lower__ZL13__ull2half_rny_nbst(shim_stream, shim_obj) -def __uint2half_rd(): +def __ull2half_rz(): pass -def _lower__ZL14__uint2half_rdj_1(shim_stream, shim_obj): +def _lower__ZL13__ull2half_rzy_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__uint2half_rdj_1(__half &retval , unsigned int* i) { - retval = __uint2half_rd(*i); + _ZL13__ull2half_rzy_nbst(__half &retval , unsigned long long* i) { + retval = __ull2half_rz(*i); return 0; } """ - _ZL14__uint2half_rdj_1 = declare_device( - "_ZL14__uint2half_rdj_1", _type___half(CPointer(uint32)) + _ZL13__ull2half_rzy_nbst = declare_device( + "_ZL13__ull2half_rzy_nbst", _type___half(CPointer(uint64)) ) - def _ZL14__uint2half_rdj_1_caller(arg_0): - return _ZL14__uint2half_rdj_1(arg_0) + def _ZL13__ull2half_rzy_nbst_caller(arg_0): + return _ZL13__ull2half_rzy_nbst(arg_0) - @lower(__uint2half_rd, uint32) + @lower(__ull2half_rz, uint64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__uint2half_rdj_1", shim_raw_str) + shim_stream.write_with_key("_ZL13__ull2half_rzy_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__uint2half_rdj_1_caller, - signature(_type___half, CPointer(uint32)), + _ZL13__ull2half_rzy_nbst_caller, + signature(_type___half, CPointer(uint64)), ptrs, ) -_lower__ZL14__uint2half_rdj_1(shim_stream, shim_obj) +_lower__ZL13__ull2half_rzy_nbst(shim_stream, shim_obj) -def __uint2half_ru(): +def __ull2half_rd(): pass -def _lower__ZL14__uint2half_ruj_1(shim_stream, shim_obj): +def _lower__ZL13__ull2half_rdy_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__uint2half_ruj_1(__half &retval , unsigned int* i) { - retval = __uint2half_ru(*i); + _ZL13__ull2half_rdy_nbst(__half &retval , unsigned long long* i) { + retval = __ull2half_rd(*i); return 0; } """ - _ZL14__uint2half_ruj_1 = declare_device( - "_ZL14__uint2half_ruj_1", _type___half(CPointer(uint32)) + _ZL13__ull2half_rdy_nbst = declare_device( + "_ZL13__ull2half_rdy_nbst", _type___half(CPointer(uint64)) ) - def _ZL14__uint2half_ruj_1_caller(arg_0): - return _ZL14__uint2half_ruj_1(arg_0) + def _ZL13__ull2half_rdy_nbst_caller(arg_0): + return _ZL13__ull2half_rdy_nbst(arg_0) - @lower(__uint2half_ru, uint32) + @lower(__ull2half_rd, uint64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL14__uint2half_ruj_1", shim_raw_str) + shim_stream.write_with_key("_ZL13__ull2half_rdy_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__uint2half_ruj_1_caller, - signature(_type___half, CPointer(uint32)), + _ZL13__ull2half_rdy_nbst_caller, + signature(_type___half, CPointer(uint64)), ptrs, ) -_lower__ZL14__uint2half_ruj_1(shim_stream, shim_obj) +_lower__ZL13__ull2half_rdy_nbst(shim_stream, shim_obj) -def __half2ushort_rn(): +def __ull2half_ru(): pass -def _lower__ZL16__half2ushort_rn6__half_1(shim_stream, shim_obj): +def _lower__ZL13__ull2half_ruy_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__half2ushort_rn6__half_1(unsigned short &retval , __half* h) { - retval = __half2ushort_rn(*h); + _ZL13__ull2half_ruy_nbst(__half &retval , unsigned long long* i) { + retval = __ull2half_ru(*i); return 0; } """ - _ZL16__half2ushort_rn6__half_1 = declare_device( - "_ZL16__half2ushort_rn6__half_1", uint16(CPointer(_type___half)) + _ZL13__ull2half_ruy_nbst = declare_device( + "_ZL13__ull2half_ruy_nbst", _type___half(CPointer(uint64)) ) - def _ZL16__half2ushort_rn6__half_1_caller(arg_0): - return _ZL16__half2ushort_rn6__half_1(arg_0) + def _ZL13__ull2half_ruy_nbst_caller(arg_0): + return _ZL13__ull2half_ruy_nbst(arg_0) - @lower(__half2ushort_rn, _type___half) + @lower(__ull2half_ru, uint64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL16__half2ushort_rn6__half_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL13__ull2half_ruy_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__half2ushort_rn6__half_1_caller, - signature(uint16, CPointer(_type___half)), + _ZL13__ull2half_ruy_nbst_caller, + signature(_type___half, CPointer(uint64)), ptrs, ) -_lower__ZL16__half2ushort_rn6__half_1(shim_stream, shim_obj) +_lower__ZL13__ull2half_ruy_nbst(shim_stream, shim_obj) -def __half2ushort_rd(): +def __half2ll_rn(): pass -def _lower__ZL16__half2ushort_rd6__half_1(shim_stream, shim_obj): +def _lower__ZL12__half2ll_rn6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__half2ushort_rd6__half_1(unsigned short &retval , __half* h) { - retval = __half2ushort_rd(*h); + _ZL12__half2ll_rn6__half_nbst(long long &retval , __half* h) { + retval = __half2ll_rn(*h); return 0; } """ - _ZL16__half2ushort_rd6__half_1 = declare_device( - "_ZL16__half2ushort_rd6__half_1", uint16(CPointer(_type___half)) + _ZL12__half2ll_rn6__half_nbst = declare_device( + "_ZL12__half2ll_rn6__half_nbst", int64(CPointer(_type___half)) ) - def _ZL16__half2ushort_rd6__half_1_caller(arg_0): - return _ZL16__half2ushort_rd6__half_1(arg_0) + def _ZL12__half2ll_rn6__half_nbst_caller(arg_0): + return _ZL12__half2ll_rn6__half_nbst(arg_0) - @lower(__half2ushort_rd, _type___half) + @lower(__half2ll_rn, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) shim_stream.write_with_key( - "_ZL16__half2ushort_rd6__half_1", shim_raw_str + "_ZL12__half2ll_rn6__half_nbst", shim_raw_str ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): @@ -2985,40 +2949,40 @@ def impl(context, builder, sig, args): return context.compile_internal( builder, - _ZL16__half2ushort_rd6__half_1_caller, - signature(uint16, CPointer(_type___half)), + _ZL12__half2ll_rn6__half_nbst_caller, + signature(int64, CPointer(_type___half)), ptrs, ) -_lower__ZL16__half2ushort_rd6__half_1(shim_stream, shim_obj) +_lower__ZL12__half2ll_rn6__half_nbst(shim_stream, shim_obj) -def __half2ushort_ru(): +def __half2ll_rd(): pass -def _lower__ZL16__half2ushort_ru6__half_1(shim_stream, shim_obj): +def _lower__ZL12__half2ll_rd6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__half2ushort_ru6__half_1(unsigned short &retval , __half* h) { - retval = __half2ushort_ru(*h); + _ZL12__half2ll_rd6__half_nbst(long long &retval , __half* h) { + retval = __half2ll_rd(*h); return 0; } """ - _ZL16__half2ushort_ru6__half_1 = declare_device( - "_ZL16__half2ushort_ru6__half_1", uint16(CPointer(_type___half)) + _ZL12__half2ll_rd6__half_nbst = declare_device( + "_ZL12__half2ll_rd6__half_nbst", int64(CPointer(_type___half)) ) - def _ZL16__half2ushort_ru6__half_1_caller(arg_0): - return _ZL16__half2ushort_ru6__half_1(arg_0) + def _ZL12__half2ll_rd6__half_nbst_caller(arg_0): + return _ZL12__half2ll_rd6__half_nbst(arg_0) - @lower(__half2ushort_ru, _type___half) + @lower(__half2ll_rd, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) shim_stream.write_with_key( - "_ZL16__half2ushort_ru6__half_1", shim_raw_str + "_ZL12__half2ll_rd6__half_nbst", shim_raw_str ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): @@ -3026,1629 +2990,1727 @@ def impl(context, builder, sig, args): return context.compile_internal( builder, - _ZL16__half2ushort_ru6__half_1_caller, - signature(uint16, CPointer(_type___half)), + _ZL12__half2ll_rd6__half_nbst_caller, + signature(int64, CPointer(_type___half)), ptrs, ) -_lower__ZL16__half2ushort_ru6__half_1(shim_stream, shim_obj) +_lower__ZL12__half2ll_rd6__half_nbst(shim_stream, shim_obj) -def __ushort2half_rn(): +def __half2ll_ru(): pass -def _lower__ZL16__ushort2half_rnt_1(shim_stream, shim_obj): +def _lower__ZL12__half2ll_ru6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__ushort2half_rnt_1(__half &retval , unsigned short* i) { - retval = __ushort2half_rn(*i); + _ZL12__half2ll_ru6__half_nbst(long long &retval , __half* h) { + retval = __half2ll_ru(*h); return 0; } """ - _ZL16__ushort2half_rnt_1 = declare_device( - "_ZL16__ushort2half_rnt_1", _type___half(CPointer(uint16)) + _ZL12__half2ll_ru6__half_nbst = declare_device( + "_ZL12__half2ll_ru6__half_nbst", int64(CPointer(_type___half)) ) - def _ZL16__ushort2half_rnt_1_caller(arg_0): - return _ZL16__ushort2half_rnt_1(arg_0) + def _ZL12__half2ll_ru6__half_nbst_caller(arg_0): + return _ZL12__half2ll_ru6__half_nbst(arg_0) - @lower(__ushort2half_rn, uint16) + @lower(__half2ll_ru, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL16__ushort2half_rnt_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL12__half2ll_ru6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__ushort2half_rnt_1_caller, - signature(_type___half, CPointer(uint16)), + _ZL12__half2ll_ru6__half_nbst_caller, + signature(int64, CPointer(_type___half)), ptrs, ) -_lower__ZL16__ushort2half_rnt_1(shim_stream, shim_obj) +_lower__ZL12__half2ll_ru6__half_nbst(shim_stream, shim_obj) -def __ushort2half_rz(): +def __ll2half_rn(): pass -def _lower__ZL16__ushort2half_rzt_1(shim_stream, shim_obj): +def _lower__ZL12__ll2half_rnx_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__ushort2half_rzt_1(__half &retval , unsigned short* i) { - retval = __ushort2half_rz(*i); + _ZL12__ll2half_rnx_nbst(__half &retval , long long* i) { + retval = __ll2half_rn(*i); return 0; } """ - _ZL16__ushort2half_rzt_1 = declare_device( - "_ZL16__ushort2half_rzt_1", _type___half(CPointer(uint16)) + _ZL12__ll2half_rnx_nbst = declare_device( + "_ZL12__ll2half_rnx_nbst", _type___half(CPointer(int64)) ) - def _ZL16__ushort2half_rzt_1_caller(arg_0): - return _ZL16__ushort2half_rzt_1(arg_0) + def _ZL12__ll2half_rnx_nbst_caller(arg_0): + return _ZL12__ll2half_rnx_nbst(arg_0) - @lower(__ushort2half_rz, uint16) + @lower(__ll2half_rn, int64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL16__ushort2half_rzt_1", shim_raw_str) + shim_stream.write_with_key("_ZL12__ll2half_rnx_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__ushort2half_rzt_1_caller, - signature(_type___half, CPointer(uint16)), + _ZL12__ll2half_rnx_nbst_caller, + signature(_type___half, CPointer(int64)), ptrs, ) -_lower__ZL16__ushort2half_rzt_1(shim_stream, shim_obj) +_lower__ZL12__ll2half_rnx_nbst(shim_stream, shim_obj) -def __ushort2half_rd(): +def __ll2half_rz(): pass -def _lower__ZL16__ushort2half_rdt_1(shim_stream, shim_obj): +def _lower__ZL12__ll2half_rzx_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__ushort2half_rdt_1(__half &retval , unsigned short* i) { - retval = __ushort2half_rd(*i); + _ZL12__ll2half_rzx_nbst(__half &retval , long long* i) { + retval = __ll2half_rz(*i); return 0; } """ - _ZL16__ushort2half_rdt_1 = declare_device( - "_ZL16__ushort2half_rdt_1", _type___half(CPointer(uint16)) + _ZL12__ll2half_rzx_nbst = declare_device( + "_ZL12__ll2half_rzx_nbst", _type___half(CPointer(int64)) ) - def _ZL16__ushort2half_rdt_1_caller(arg_0): - return _ZL16__ushort2half_rdt_1(arg_0) + def _ZL12__ll2half_rzx_nbst_caller(arg_0): + return _ZL12__ll2half_rzx_nbst(arg_0) - @lower(__ushort2half_rd, uint16) + @lower(__ll2half_rz, int64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL16__ushort2half_rdt_1", shim_raw_str) + shim_stream.write_with_key("_ZL12__ll2half_rzx_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__ushort2half_rdt_1_caller, - signature(_type___half, CPointer(uint16)), + _ZL12__ll2half_rzx_nbst_caller, + signature(_type___half, CPointer(int64)), ptrs, ) -_lower__ZL16__ushort2half_rdt_1(shim_stream, shim_obj) +_lower__ZL12__ll2half_rzx_nbst(shim_stream, shim_obj) -def __ushort2half_ru(): +def __ll2half_rd(): pass -def _lower__ZL16__ushort2half_rut_1(shim_stream, shim_obj): +def _lower__ZL12__ll2half_rdx_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__ushort2half_rut_1(__half &retval , unsigned short* i) { - retval = __ushort2half_ru(*i); + _ZL12__ll2half_rdx_nbst(__half &retval , long long* i) { + retval = __ll2half_rd(*i); return 0; } """ - _ZL16__ushort2half_rut_1 = declare_device( - "_ZL16__ushort2half_rut_1", _type___half(CPointer(uint16)) + _ZL12__ll2half_rdx_nbst = declare_device( + "_ZL12__ll2half_rdx_nbst", _type___half(CPointer(int64)) ) - def _ZL16__ushort2half_rut_1_caller(arg_0): - return _ZL16__ushort2half_rut_1(arg_0) + def _ZL12__ll2half_rdx_nbst_caller(arg_0): + return _ZL12__ll2half_rdx_nbst(arg_0) - @lower(__ushort2half_ru, uint16) + @lower(__ll2half_rd, int64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL16__ushort2half_rut_1", shim_raw_str) + shim_stream.write_with_key("_ZL12__ll2half_rdx_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__ushort2half_rut_1_caller, - signature(_type___half, CPointer(uint16)), + _ZL12__ll2half_rdx_nbst_caller, + signature(_type___half, CPointer(int64)), ptrs, ) -_lower__ZL16__ushort2half_rut_1(shim_stream, shim_obj) +_lower__ZL12__ll2half_rdx_nbst(shim_stream, shim_obj) -def __half2ull_rn(): +def __ll2half_ru(): pass -def _lower__ZL13__half2ull_rn6__half_1(shim_stream, shim_obj): +def _lower__ZL12__ll2half_rux_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2ull_rn6__half_1(unsigned long long &retval , __half* h) { - retval = __half2ull_rn(*h); + _ZL12__ll2half_rux_nbst(__half &retval , long long* i) { + retval = __ll2half_ru(*i); return 0; } """ - _ZL13__half2ull_rn6__half_1 = declare_device( - "_ZL13__half2ull_rn6__half_1", uint64(CPointer(_type___half)) + _ZL12__ll2half_rux_nbst = declare_device( + "_ZL12__ll2half_rux_nbst", _type___half(CPointer(int64)) ) - def _ZL13__half2ull_rn6__half_1_caller(arg_0): - return _ZL13__half2ull_rn6__half_1(arg_0) + def _ZL12__ll2half_rux_nbst_caller(arg_0): + return _ZL12__ll2half_rux_nbst(arg_0) - @lower(__half2ull_rn, _type___half) + @lower(__ll2half_ru, int64) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2ull_rn6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL12__ll2half_rux_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2ull_rn6__half_1_caller, - signature(uint64, CPointer(_type___half)), + _ZL12__ll2half_rux_nbst_caller, + signature(_type___half, CPointer(int64)), ptrs, ) -_lower__ZL13__half2ull_rn6__half_1(shim_stream, shim_obj) +_lower__ZL12__ll2half_rux_nbst(shim_stream, shim_obj) -def __half2ull_rd(): +def htrunc(): pass -def _lower__ZL13__half2ull_rd6__half_1(shim_stream, shim_obj): +def _lower__ZL6htrunc6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2ull_rd6__half_1(unsigned long long &retval , __half* h) { - retval = __half2ull_rd(*h); + _ZL6htrunc6__half_nbst(__half &retval , __half* h) { + retval = htrunc(*h); return 0; } """ - _ZL13__half2ull_rd6__half_1 = declare_device( - "_ZL13__half2ull_rd6__half_1", uint64(CPointer(_type___half)) + _ZL6htrunc6__half_nbst = declare_device( + "_ZL6htrunc6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL13__half2ull_rd6__half_1_caller(arg_0): - return _ZL13__half2ull_rd6__half_1(arg_0) + def _ZL6htrunc6__half_nbst_caller(arg_0): + return _ZL6htrunc6__half_nbst(arg_0) - @lower(__half2ull_rd, _type___half) + @lower(htrunc, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2ull_rd6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL6htrunc6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2ull_rd6__half_1_caller, - signature(uint64, CPointer(_type___half)), + _ZL6htrunc6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL13__half2ull_rd6__half_1(shim_stream, shim_obj) +_lower__ZL6htrunc6__half_nbst(shim_stream, shim_obj) -def __half2ull_ru(): +def hceil(): pass -def _lower__ZL13__half2ull_ru6__half_1(shim_stream, shim_obj): +def _lower__ZL5hceil6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__half2ull_ru6__half_1(unsigned long long &retval , __half* h) { - retval = __half2ull_ru(*h); + _ZL5hceil6__half_nbst(__half &retval , __half* h) { + retval = hceil(*h); return 0; } """ - _ZL13__half2ull_ru6__half_1 = declare_device( - "_ZL13__half2ull_ru6__half_1", uint64(CPointer(_type___half)) + _ZL5hceil6__half_nbst = declare_device( + "_ZL5hceil6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL13__half2ull_ru6__half_1_caller(arg_0): - return _ZL13__half2ull_ru6__half_1(arg_0) + def _ZL5hceil6__half_nbst_caller(arg_0): + return _ZL5hceil6__half_nbst(arg_0) - @lower(__half2ull_ru, _type___half) + @lower(hceil, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__half2ull_ru6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL5hceil6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__half2ull_ru6__half_1_caller, - signature(uint64, CPointer(_type___half)), + _ZL5hceil6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL13__half2ull_ru6__half_1(shim_stream, shim_obj) +_lower__ZL5hceil6__half_nbst(shim_stream, shim_obj) -def __ull2half_rn(): +def hfloor(): pass -def _lower__ZL13__ull2half_rny_1(shim_stream, shim_obj): +def _lower__ZL6hfloor6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__ull2half_rny_1(__half &retval , unsigned long long* i) { - retval = __ull2half_rn(*i); + _ZL6hfloor6__half_nbst(__half &retval , __half* h) { + retval = hfloor(*h); return 0; } """ - _ZL13__ull2half_rny_1 = declare_device( - "_ZL13__ull2half_rny_1", _type___half(CPointer(uint64)) + _ZL6hfloor6__half_nbst = declare_device( + "_ZL6hfloor6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL13__ull2half_rny_1_caller(arg_0): - return _ZL13__ull2half_rny_1(arg_0) + def _ZL6hfloor6__half_nbst_caller(arg_0): + return _ZL6hfloor6__half_nbst(arg_0) - @lower(__ull2half_rn, uint64) + @lower(hfloor, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__ull2half_rny_1", shim_raw_str) + shim_stream.write_with_key("_ZL6hfloor6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__ull2half_rny_1_caller, - signature(_type___half, CPointer(uint64)), + _ZL6hfloor6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL13__ull2half_rny_1(shim_stream, shim_obj) +_lower__ZL6hfloor6__half_nbst(shim_stream, shim_obj) -def __ull2half_rz(): +def hrint(): pass -def _lower__ZL13__ull2half_rzy_1(shim_stream, shim_obj): +def _lower__ZL5hrint6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__ull2half_rzy_1(__half &retval , unsigned long long* i) { - retval = __ull2half_rz(*i); + _ZL5hrint6__half_nbst(__half &retval , __half* h) { + retval = hrint(*h); return 0; } """ - _ZL13__ull2half_rzy_1 = declare_device( - "_ZL13__ull2half_rzy_1", _type___half(CPointer(uint64)) + _ZL5hrint6__half_nbst = declare_device( + "_ZL5hrint6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL13__ull2half_rzy_1_caller(arg_0): - return _ZL13__ull2half_rzy_1(arg_0) + def _ZL5hrint6__half_nbst_caller(arg_0): + return _ZL5hrint6__half_nbst(arg_0) - @lower(__ull2half_rz, uint64) + @lower(hrint, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__ull2half_rzy_1", shim_raw_str) + shim_stream.write_with_key("_ZL5hrint6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__ull2half_rzy_1_caller, - signature(_type___half, CPointer(uint64)), + _ZL5hrint6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL13__ull2half_rzy_1(shim_stream, shim_obj) +_lower__ZL5hrint6__half_nbst(shim_stream, shim_obj) -def __ull2half_rd(): +def __hisinf(): pass -def _lower__ZL13__ull2half_rdy_1(shim_stream, shim_obj): +def _lower__ZL8__hisinf6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__ull2half_rdy_1(__half &retval , unsigned long long* i) { - retval = __ull2half_rd(*i); + _ZL8__hisinf6__half_nbst(int &retval , __half* a) { + retval = __hisinf(*a); return 0; } """ - _ZL13__ull2half_rdy_1 = declare_device( - "_ZL13__ull2half_rdy_1", _type___half(CPointer(uint64)) + _ZL8__hisinf6__half_nbst = declare_device( + "_ZL8__hisinf6__half_nbst", int32(CPointer(_type___half)) ) - def _ZL13__ull2half_rdy_1_caller(arg_0): - return _ZL13__ull2half_rdy_1(arg_0) + def _ZL8__hisinf6__half_nbst_caller(arg_0): + return _ZL8__hisinf6__half_nbst(arg_0) - @lower(__ull2half_rd, uint64) + @lower(__hisinf, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__ull2half_rdy_1", shim_raw_str) + shim_stream.write_with_key("_ZL8__hisinf6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__ull2half_rdy_1_caller, - signature(_type___half, CPointer(uint64)), + _ZL8__hisinf6__half_nbst_caller, + signature(int32, CPointer(_type___half)), ptrs, ) -_lower__ZL13__ull2half_rdy_1(shim_stream, shim_obj) +_lower__ZL8__hisinf6__half_nbst(shim_stream, shim_obj) -def __ull2half_ru(): +def __half_as_short(): pass -def _lower__ZL13__ull2half_ruy_1(shim_stream, shim_obj): +def _lower__ZL15__half_as_short6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__ull2half_ruy_1(__half &retval , unsigned long long* i) { - retval = __ull2half_ru(*i); + _ZL15__half_as_short6__half_nbst(short &retval , __half* h) { + retval = __half_as_short(*h); return 0; } """ - _ZL13__ull2half_ruy_1 = declare_device( - "_ZL13__ull2half_ruy_1", _type___half(CPointer(uint64)) + _ZL15__half_as_short6__half_nbst = declare_device( + "_ZL15__half_as_short6__half_nbst", int16(CPointer(_type___half)) ) - def _ZL13__ull2half_ruy_1_caller(arg_0): - return _ZL13__ull2half_ruy_1(arg_0) + def _ZL15__half_as_short6__half_nbst_caller(arg_0): + return _ZL15__half_as_short6__half_nbst(arg_0) - @lower(__ull2half_ru, uint64) + @lower(__half_as_short, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13__ull2half_ruy_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL15__half_as_short6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__ull2half_ruy_1_caller, - signature(_type___half, CPointer(uint64)), + _ZL15__half_as_short6__half_nbst_caller, + signature(int16, CPointer(_type___half)), ptrs, ) -_lower__ZL13__ull2half_ruy_1(shim_stream, shim_obj) +_lower__ZL15__half_as_short6__half_nbst(shim_stream, shim_obj) -def __half2ll_rn(): +def __half_as_ushort(): pass -def _lower__ZL12__half2ll_rn6__half_1(shim_stream, shim_obj): +def _lower__ZL16__half_as_ushort6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__half2ll_rn6__half_1(long long &retval , __half* h) { - retval = __half2ll_rn(*h); + _ZL16__half_as_ushort6__half_nbst(unsigned short &retval , __half* h) { + retval = __half_as_ushort(*h); return 0; } """ - _ZL12__half2ll_rn6__half_1 = declare_device( - "_ZL12__half2ll_rn6__half_1", int64(CPointer(_type___half)) + _ZL16__half_as_ushort6__half_nbst = declare_device( + "_ZL16__half_as_ushort6__half_nbst", uint16(CPointer(_type___half)) ) - def _ZL12__half2ll_rn6__half_1_caller(arg_0): - return _ZL12__half2ll_rn6__half_1(arg_0) + def _ZL16__half_as_ushort6__half_nbst_caller(arg_0): + return _ZL16__half_as_ushort6__half_nbst(arg_0) - @lower(__half2ll_rn, _type___half) + @lower(__half_as_ushort, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__half2ll_rn6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL16__half_as_ushort6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__half2ll_rn6__half_1_caller, - signature(int64, CPointer(_type___half)), + _ZL16__half_as_ushort6__half_nbst_caller, + signature(uint16, CPointer(_type___half)), ptrs, ) -_lower__ZL12__half2ll_rn6__half_1(shim_stream, shim_obj) +_lower__ZL16__half_as_ushort6__half_nbst(shim_stream, shim_obj) -def __half2ll_rd(): +def __short_as_half(): pass -def _lower__ZL12__half2ll_rd6__half_1(shim_stream, shim_obj): +def _lower__ZL15__short_as_halfs_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__half2ll_rd6__half_1(long long &retval , __half* h) { - retval = __half2ll_rd(*h); + _ZL15__short_as_halfs_nbst(__half &retval , short* i) { + retval = __short_as_half(*i); return 0; } """ - _ZL12__half2ll_rd6__half_1 = declare_device( - "_ZL12__half2ll_rd6__half_1", int64(CPointer(_type___half)) + _ZL15__short_as_halfs_nbst = declare_device( + "_ZL15__short_as_halfs_nbst", _type___half(CPointer(int16)) ) - def _ZL12__half2ll_rd6__half_1_caller(arg_0): - return _ZL12__half2ll_rd6__half_1(arg_0) + def _ZL15__short_as_halfs_nbst_caller(arg_0): + return _ZL15__short_as_halfs_nbst(arg_0) - @lower(__half2ll_rd, _type___half) + @lower(__short_as_half, int16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__half2ll_rd6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL15__short_as_halfs_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__half2ll_rd6__half_1_caller, - signature(int64, CPointer(_type___half)), + _ZL15__short_as_halfs_nbst_caller, + signature(_type___half, CPointer(int16)), ptrs, ) -_lower__ZL12__half2ll_rd6__half_1(shim_stream, shim_obj) +_lower__ZL15__short_as_halfs_nbst(shim_stream, shim_obj) -def __half2ll_ru(): +def __ushort_as_half(): pass -def _lower__ZL12__half2ll_ru6__half_1(shim_stream, shim_obj): +def _lower__ZL16__ushort_as_halft_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__half2ll_ru6__half_1(long long &retval , __half* h) { - retval = __half2ll_ru(*h); + _ZL16__ushort_as_halft_nbst(__half &retval , unsigned short* i) { + retval = __ushort_as_half(*i); return 0; } """ - _ZL12__half2ll_ru6__half_1 = declare_device( - "_ZL12__half2ll_ru6__half_1", int64(CPointer(_type___half)) + _ZL16__ushort_as_halft_nbst = declare_device( + "_ZL16__ushort_as_halft_nbst", _type___half(CPointer(uint16)) ) - def _ZL12__half2ll_ru6__half_1_caller(arg_0): - return _ZL12__half2ll_ru6__half_1(arg_0) + def _ZL16__ushort_as_halft_nbst_caller(arg_0): + return _ZL16__ushort_as_halft_nbst(arg_0) - @lower(__half2ll_ru, _type___half) + @lower(__ushort_as_half, uint16) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__half2ll_ru6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL16__ushort_as_halft_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__half2ll_ru6__half_1_caller, - signature(int64, CPointer(_type___half)), + _ZL16__ushort_as_halft_nbst_caller, + signature(_type___half, CPointer(uint16)), ptrs, ) -_lower__ZL12__half2ll_ru6__half_1(shim_stream, shim_obj) +_lower__ZL16__ushort_as_halft_nbst(shim_stream, shim_obj) -def __ll2half_rn(): +def __hmax(): pass -def _lower__ZL12__ll2half_rnx_1(shim_stream, shim_obj): +def _lower__ZL6__hmax6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__ll2half_rnx_1(__half &retval , long long* i) { - retval = __ll2half_rn(*i); + _ZL6__hmax6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hmax(*a, *b); return 0; } """ - _ZL12__ll2half_rnx_1 = declare_device( - "_ZL12__ll2half_rnx_1", _type___half(CPointer(int64)) + _ZL6__hmax6__halfS__nbst = declare_device( + "_ZL6__hmax6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL12__ll2half_rnx_1_caller(arg_0): - return _ZL12__ll2half_rnx_1(arg_0) + def _ZL6__hmax6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hmax6__halfS__nbst(arg_0, arg_1) - @lower(__ll2half_rn, int64) + @lower(__hmax, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__ll2half_rnx_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hmax6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__ll2half_rnx_1_caller, - signature(_type___half, CPointer(int64)), + _ZL6__hmax6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL12__ll2half_rnx_1(shim_stream, shim_obj) +_lower__ZL6__hmax6__halfS__nbst(shim_stream, shim_obj) -def __ll2half_rz(): +def __hmin(): pass -def _lower__ZL12__ll2half_rzx_1(shim_stream, shim_obj): +def _lower__ZL6__hmin6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__ll2half_rzx_1(__half &retval , long long* i) { - retval = __ll2half_rz(*i); + _ZL6__hmin6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hmin(*a, *b); return 0; } """ - _ZL12__ll2half_rzx_1 = declare_device( - "_ZL12__ll2half_rzx_1", _type___half(CPointer(int64)) + _ZL6__hmin6__halfS__nbst = declare_device( + "_ZL6__hmin6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL12__ll2half_rzx_1_caller(arg_0): - return _ZL12__ll2half_rzx_1(arg_0) + def _ZL6__hmin6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hmin6__halfS__nbst(arg_0, arg_1) - @lower(__ll2half_rz, int64) + @lower(__hmin, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__ll2half_rzx_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hmin6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__ll2half_rzx_1_caller, - signature(_type___half, CPointer(int64)), + _ZL6__hmin6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL12__ll2half_rzx_1(shim_stream, shim_obj) +_lower__ZL6__hmin6__halfS__nbst(shim_stream, shim_obj) -def __ll2half_rd(): +def __shfl_sync(): pass -def _lower__ZL12__ll2half_rdx_1(shim_stream, shim_obj): +def _lower__ZL11__shfl_syncj6__halfii_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__ll2half_rdx_1(__half &retval , long long* i) { - retval = __ll2half_rd(*i); + _ZL11__shfl_syncj6__halfii_nbst(__half &retval , unsigned int* mask, __half* var, int* srcLane, int* width) { + retval = __shfl_sync(*mask, *var, *srcLane, *width); return 0; } """ - _ZL12__ll2half_rdx_1 = declare_device( - "_ZL12__ll2half_rdx_1", _type___half(CPointer(int64)) + _ZL11__shfl_syncj6__halfii_nbst = declare_device( + "_ZL11__shfl_syncj6__halfii_nbst", + _type___half( + CPointer(uint32), + CPointer(_type___half), + CPointer(int32), + CPointer(int32), + ), ) - def _ZL12__ll2half_rdx_1_caller(arg_0): - return _ZL12__ll2half_rdx_1(arg_0) + def _ZL11__shfl_syncj6__halfii_nbst_caller(arg_0, arg_1, arg_2, arg_3): + return _ZL11__shfl_syncj6__halfii_nbst(arg_0, arg_1, arg_2, arg_3) - @lower(__ll2half_rd, int64) + @lower(__shfl_sync, uint32, _type___half, int32, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__ll2half_rdx_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL11__shfl_syncj6__halfii_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__ll2half_rdx_1_caller, - signature(_type___half, CPointer(int64)), + _ZL11__shfl_syncj6__halfii_nbst_caller, + signature( + _type___half, + CPointer(uint32), + CPointer(_type___half), + CPointer(int32), + CPointer(int32), + ), ptrs, ) -_lower__ZL12__ll2half_rdx_1(shim_stream, shim_obj) +_lower__ZL11__shfl_syncj6__halfii_nbst(shim_stream, shim_obj) -def __ll2half_ru(): +def __shfl_up_sync(): pass -def _lower__ZL12__ll2half_rux_1(shim_stream, shim_obj): +def _lower__ZL14__shfl_up_syncj6__halfji_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__ll2half_rux_1(__half &retval , long long* i) { - retval = __ll2half_ru(*i); + _ZL14__shfl_up_syncj6__halfji_nbst(__half &retval , unsigned int* mask, __half* var, unsigned int* delta, int* width) { + retval = __shfl_up_sync(*mask, *var, *delta, *width); return 0; } """ - _ZL12__ll2half_rux_1 = declare_device( - "_ZL12__ll2half_rux_1", _type___half(CPointer(int64)) + _ZL14__shfl_up_syncj6__halfji_nbst = declare_device( + "_ZL14__shfl_up_syncj6__halfji_nbst", + _type___half( + CPointer(uint32), + CPointer(_type___half), + CPointer(uint32), + CPointer(int32), + ), ) - def _ZL12__ll2half_rux_1_caller(arg_0): - return _ZL12__ll2half_rux_1(arg_0) + def _ZL14__shfl_up_syncj6__halfji_nbst_caller(arg_0, arg_1, arg_2, arg_3): + return _ZL14__shfl_up_syncj6__halfji_nbst(arg_0, arg_1, arg_2, arg_3) - @lower(__ll2half_ru, int64) + @lower(__shfl_up_sync, uint32, _type___half, uint32, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__ll2half_rux_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL14__shfl_up_syncj6__halfji_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__ll2half_rux_1_caller, - signature(_type___half, CPointer(int64)), + _ZL14__shfl_up_syncj6__halfji_nbst_caller, + signature( + _type___half, + CPointer(uint32), + CPointer(_type___half), + CPointer(uint32), + CPointer(int32), + ), ptrs, ) -_lower__ZL12__ll2half_rux_1(shim_stream, shim_obj) +_lower__ZL14__shfl_up_syncj6__halfji_nbst(shim_stream, shim_obj) -def htrunc(): +def __shfl_down_sync(): pass -def _lower__ZL6htrunc6__half_1(shim_stream, shim_obj): +def _lower__ZL16__shfl_down_syncj6__halfji_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6htrunc6__half_1(__half &retval , __half* h) { - retval = htrunc(*h); + _ZL16__shfl_down_syncj6__halfji_nbst(__half &retval , unsigned int* mask, __half* var, unsigned int* delta, int* width) { + retval = __shfl_down_sync(*mask, *var, *delta, *width); return 0; } """ - _ZL6htrunc6__half_1 = declare_device( - "_ZL6htrunc6__half_1", _type___half(CPointer(_type___half)) + _ZL16__shfl_down_syncj6__halfji_nbst = declare_device( + "_ZL16__shfl_down_syncj6__halfji_nbst", + _type___half( + CPointer(uint32), + CPointer(_type___half), + CPointer(uint32), + CPointer(int32), + ), ) - def _ZL6htrunc6__half_1_caller(arg_0): - return _ZL6htrunc6__half_1(arg_0) + def _ZL16__shfl_down_syncj6__halfji_nbst_caller(arg_0, arg_1, arg_2, arg_3): + return _ZL16__shfl_down_syncj6__halfji_nbst(arg_0, arg_1, arg_2, arg_3) - @lower(htrunc, _type___half) + @lower(__shfl_down_sync, uint32, _type___half, uint32, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6htrunc6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL16__shfl_down_syncj6__halfji_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6htrunc6__half_1_caller, - signature(_type___half, CPointer(_type___half)), + _ZL16__shfl_down_syncj6__halfji_nbst_caller, + signature( + _type___half, + CPointer(uint32), + CPointer(_type___half), + CPointer(uint32), + CPointer(int32), + ), ptrs, ) -_lower__ZL6htrunc6__half_1(shim_stream, shim_obj) +_lower__ZL16__shfl_down_syncj6__halfji_nbst(shim_stream, shim_obj) -def hceil(): +def __shfl_xor_sync(): pass -def _lower__ZL5hceil6__half_1(shim_stream, shim_obj): +def _lower__ZL15__shfl_xor_syncj6__halfii_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL5hceil6__half_1(__half &retval , __half* h) { - retval = hceil(*h); + _ZL15__shfl_xor_syncj6__halfii_nbst(__half &retval , unsigned int* mask, __half* var, int* laneMask, int* width) { + retval = __shfl_xor_sync(*mask, *var, *laneMask, *width); return 0; } """ - _ZL5hceil6__half_1 = declare_device( - "_ZL5hceil6__half_1", _type___half(CPointer(_type___half)) + _ZL15__shfl_xor_syncj6__halfii_nbst = declare_device( + "_ZL15__shfl_xor_syncj6__halfii_nbst", + _type___half( + CPointer(uint32), + CPointer(_type___half), + CPointer(int32), + CPointer(int32), + ), ) - def _ZL5hceil6__half_1_caller(arg_0): - return _ZL5hceil6__half_1(arg_0) + def _ZL15__shfl_xor_syncj6__halfii_nbst_caller(arg_0, arg_1, arg_2, arg_3): + return _ZL15__shfl_xor_syncj6__halfii_nbst(arg_0, arg_1, arg_2, arg_3) - @lower(hceil, _type___half) + @lower(__shfl_xor_sync, uint32, _type___half, int32, int32) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5hceil6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL15__shfl_xor_syncj6__halfii_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL5hceil6__half_1_caller, - signature(_type___half, CPointer(_type___half)), + _ZL15__shfl_xor_syncj6__halfii_nbst_caller, + signature( + _type___half, + CPointer(uint32), + CPointer(_type___half), + CPointer(int32), + CPointer(int32), + ), ptrs, ) -_lower__ZL5hceil6__half_1(shim_stream, shim_obj) +_lower__ZL15__shfl_xor_syncj6__halfii_nbst(shim_stream, shim_obj) -def hfloor(): +def __ldg(): pass -def _lower__ZL6hfloor6__half_1(shim_stream, shim_obj): +def _lower__ZL5__ldgPK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6hfloor6__half_1(__half &retval , __half* h) { - retval = hfloor(*h); + _ZL5__ldgPK6__half_nbst(__half &retval , __half ** ptr) { + retval = __ldg(*ptr); return 0; } """ - _ZL6hfloor6__half_1 = declare_device( - "_ZL6hfloor6__half_1", _type___half(CPointer(_type___half)) + _ZL5__ldgPK6__half_nbst = declare_device( + "_ZL5__ldgPK6__half_nbst", + _type___half(CPointer(CPointer(_type___half))), ) - def _ZL6hfloor6__half_1_caller(arg_0): - return _ZL6hfloor6__half_1(arg_0) + def _ZL5__ldgPK6__half_nbst_caller(arg_0): + return _ZL5__ldgPK6__half_nbst(arg_0) - @lower(hfloor, _type___half) + @lower(__ldg, CPointer(_type___half)) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6hfloor6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL5__ldgPK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6hfloor6__half_1_caller, - signature(_type___half, CPointer(_type___half)), + _ZL5__ldgPK6__half_nbst_caller, + signature(_type___half, CPointer(CPointer(_type___half))), ptrs, ) -_lower__ZL6hfloor6__half_1(shim_stream, shim_obj) +_lower__ZL5__ldgPK6__half_nbst(shim_stream, shim_obj) -def hrint(): +def __ldcg(): pass -def _lower__ZL5hrint6__half_1(shim_stream, shim_obj): +def _lower__ZL6__ldcgPK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL5hrint6__half_1(__half &retval , __half* h) { - retval = hrint(*h); + _ZL6__ldcgPK6__half_nbst(__half &retval , __half ** ptr) { + retval = __ldcg(*ptr); return 0; } """ - _ZL5hrint6__half_1 = declare_device( - "_ZL5hrint6__half_1", _type___half(CPointer(_type___half)) + _ZL6__ldcgPK6__half_nbst = declare_device( + "_ZL6__ldcgPK6__half_nbst", + _type___half(CPointer(CPointer(_type___half))), ) - def _ZL5hrint6__half_1_caller(arg_0): - return _ZL5hrint6__half_1(arg_0) + def _ZL6__ldcgPK6__half_nbst_caller(arg_0): + return _ZL6__ldcgPK6__half_nbst(arg_0) - @lower(hrint, _type___half) + @lower(__ldcg, CPointer(_type___half)) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5hrint6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__ldcgPK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL5hrint6__half_1_caller, - signature(_type___half, CPointer(_type___half)), + _ZL6__ldcgPK6__half_nbst_caller, + signature(_type___half, CPointer(CPointer(_type___half))), ptrs, ) -_lower__ZL5hrint6__half_1(shim_stream, shim_obj) +_lower__ZL6__ldcgPK6__half_nbst(shim_stream, shim_obj) -def h2trunc(): +def __ldca(): pass -def _lower__ZL7h2trunc7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__ldcaPK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7h2trunc7__half2_1(__half2 &retval , __half2* h) { - retval = h2trunc(*h); + _ZL6__ldcaPK6__half_nbst(__half &retval , __half ** ptr) { + retval = __ldca(*ptr); return 0; } """ - _ZL7h2trunc7__half2_1 = declare_device( - "_ZL7h2trunc7__half2_1", _type___half2(CPointer(_type___half2)) + _ZL6__ldcaPK6__half_nbst = declare_device( + "_ZL6__ldcaPK6__half_nbst", + _type___half(CPointer(CPointer(_type___half))), ) - def _ZL7h2trunc7__half2_1_caller(arg_0): - return _ZL7h2trunc7__half2_1(arg_0) + def _ZL6__ldcaPK6__half_nbst_caller(arg_0): + return _ZL6__ldcaPK6__half_nbst(arg_0) - @lower(h2trunc, _type___half2) + @lower(__ldca, CPointer(_type___half)) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7h2trunc7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__ldcaPK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7h2trunc7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), + _ZL6__ldcaPK6__half_nbst_caller, + signature(_type___half, CPointer(CPointer(_type___half))), ptrs, ) -_lower__ZL7h2trunc7__half2_1(shim_stream, shim_obj) +_lower__ZL6__ldcaPK6__half_nbst(shim_stream, shim_obj) -def h2ceil(): +def __ldcs(): pass -def _lower__ZL6h2ceil7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__ldcsPK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6h2ceil7__half2_1(__half2 &retval , __half2* h) { - retval = h2ceil(*h); + _ZL6__ldcsPK6__half_nbst(__half &retval , __half ** ptr) { + retval = __ldcs(*ptr); return 0; } """ - _ZL6h2ceil7__half2_1 = declare_device( - "_ZL6h2ceil7__half2_1", _type___half2(CPointer(_type___half2)) + _ZL6__ldcsPK6__half_nbst = declare_device( + "_ZL6__ldcsPK6__half_nbst", + _type___half(CPointer(CPointer(_type___half))), ) - def _ZL6h2ceil7__half2_1_caller(arg_0): - return _ZL6h2ceil7__half2_1(arg_0) + def _ZL6__ldcsPK6__half_nbst_caller(arg_0): + return _ZL6__ldcsPK6__half_nbst(arg_0) - @lower(h2ceil, _type___half2) + @lower(__ldcs, CPointer(_type___half)) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6h2ceil7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__ldcsPK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6h2ceil7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), + _ZL6__ldcsPK6__half_nbst_caller, + signature(_type___half, CPointer(CPointer(_type___half))), ptrs, ) -_lower__ZL6h2ceil7__half2_1(shim_stream, shim_obj) +_lower__ZL6__ldcsPK6__half_nbst(shim_stream, shim_obj) -def h2floor(): +def __ldlu(): pass -def _lower__ZL7h2floor7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__ldluPK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7h2floor7__half2_1(__half2 &retval , __half2* h) { - retval = h2floor(*h); + _ZL6__ldluPK6__half_nbst(__half &retval , __half ** ptr) { + retval = __ldlu(*ptr); return 0; } """ - _ZL7h2floor7__half2_1 = declare_device( - "_ZL7h2floor7__half2_1", _type___half2(CPointer(_type___half2)) + _ZL6__ldluPK6__half_nbst = declare_device( + "_ZL6__ldluPK6__half_nbst", + _type___half(CPointer(CPointer(_type___half))), ) - def _ZL7h2floor7__half2_1_caller(arg_0): - return _ZL7h2floor7__half2_1(arg_0) + def _ZL6__ldluPK6__half_nbst_caller(arg_0): + return _ZL6__ldluPK6__half_nbst(arg_0) - @lower(h2floor, _type___half2) + @lower(__ldlu, CPointer(_type___half)) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7h2floor7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__ldluPK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7h2floor7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), + _ZL6__ldluPK6__half_nbst_caller, + signature(_type___half, CPointer(CPointer(_type___half))), ptrs, ) -_lower__ZL7h2floor7__half2_1(shim_stream, shim_obj) +_lower__ZL6__ldluPK6__half_nbst(shim_stream, shim_obj) -def h2rint(): +def __ldcv(): pass -def _lower__ZL6h2rint7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__ldcvPK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6h2rint7__half2_1(__half2 &retval , __half2* h) { - retval = h2rint(*h); + _ZL6__ldcvPK6__half_nbst(__half &retval , __half ** ptr) { + retval = __ldcv(*ptr); return 0; } """ - _ZL6h2rint7__half2_1 = declare_device( - "_ZL6h2rint7__half2_1", _type___half2(CPointer(_type___half2)) + _ZL6__ldcvPK6__half_nbst = declare_device( + "_ZL6__ldcvPK6__half_nbst", + _type___half(CPointer(CPointer(_type___half))), ) - def _ZL6h2rint7__half2_1_caller(arg_0): - return _ZL6h2rint7__half2_1(arg_0) + def _ZL6__ldcvPK6__half_nbst_caller(arg_0): + return _ZL6__ldcvPK6__half_nbst(arg_0) - @lower(h2rint, _type___half2) + @lower(__ldcv, CPointer(_type___half)) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6h2rint7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__ldcvPK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6h2rint7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), + _ZL6__ldcvPK6__half_nbst_caller, + signature(_type___half, CPointer(CPointer(_type___half))), ptrs, ) -_lower__ZL6h2rint7__half2_1(shim_stream, shim_obj) +_lower__ZL6__ldcvPK6__half_nbst(shim_stream, shim_obj) -def __half2half2(): +def __stwb(): pass -def _lower__ZL12__half2half26__half_1(shim_stream, shim_obj): +def _lower__ZL6__stwbP6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__half2half26__half_1(__half2 &retval , __half* a) { - retval = __half2half2(*a); + _ZL6__stwbP6__halfS__nbst(int &retval , __half ** ptr, __half* value) { + __stwb(*ptr, *value); return 0; } """ - _ZL12__half2half26__half_1 = declare_device( - "_ZL12__half2half26__half_1", _type___half2(CPointer(_type___half)) + _ZL6__stwbP6__halfS__nbst = declare_device( + "_ZL6__stwbP6__halfS__nbst", + void(CPointer(CPointer(_type___half)), CPointer(_type___half)), ) - def _ZL12__half2half26__half_1_caller(arg_0): - return _ZL12__half2half26__half_1(arg_0) + def _ZL6__stwbP6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__stwbP6__halfS__nbst(arg_0, arg_1) - @lower(__half2half2, _type___half) + @lower(__stwb, CPointer(_type___half), _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__half2half26__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__stwbP6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__half2half26__half_1_caller, - signature(_type___half2, CPointer(_type___half)), + _ZL6__stwbP6__halfS__nbst_caller, + signature( + void, CPointer(CPointer(_type___half)), CPointer(_type___half) + ), ptrs, ) -_lower__ZL12__half2half26__half_1(shim_stream, shim_obj) +_lower__ZL6__stwbP6__halfS__nbst(shim_stream, shim_obj) -def __lowhigh2highlow(): +def __stcg(): pass -def _lower__ZL17__lowhigh2highlow7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__stcgP6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL17__lowhigh2highlow7__half2_1(__half2 &retval , __half2* a) { - retval = __lowhigh2highlow(*a); + _ZL6__stcgP6__halfS__nbst(int &retval , __half ** ptr, __half* value) { + __stcg(*ptr, *value); return 0; } """ - _ZL17__lowhigh2highlow7__half2_1 = declare_device( - "_ZL17__lowhigh2highlow7__half2_1", - _type___half2(CPointer(_type___half2)), + _ZL6__stcgP6__halfS__nbst = declare_device( + "_ZL6__stcgP6__halfS__nbst", + void(CPointer(CPointer(_type___half)), CPointer(_type___half)), ) - def _ZL17__lowhigh2highlow7__half2_1_caller(arg_0): - return _ZL17__lowhigh2highlow7__half2_1(arg_0) + def _ZL6__stcgP6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__stcgP6__halfS__nbst(arg_0, arg_1) - @lower(__lowhigh2highlow, _type___half2) + @lower(__stcg, CPointer(_type___half), _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL17__lowhigh2highlow7__half2_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL6__stcgP6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL17__lowhigh2highlow7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), + _ZL6__stcgP6__halfS__nbst_caller, + signature( + void, CPointer(CPointer(_type___half)), CPointer(_type___half) + ), ptrs, ) -_lower__ZL17__lowhigh2highlow7__half2_1(shim_stream, shim_obj) +_lower__ZL6__stcgP6__halfS__nbst(shim_stream, shim_obj) -def __lows2half2(): +def __stcs(): pass -def _lower__ZL12__lows2half27__half2S__1(shim_stream, shim_obj): +def _lower__ZL6__stcsP6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__lows2half27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __lows2half2(*a, *b); + _ZL6__stcsP6__halfS__nbst(int &retval , __half ** ptr, __half* value) { + __stcs(*ptr, *value); return 0; } """ - _ZL12__lows2half27__half2S__1 = declare_device( - "_ZL12__lows2half27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZL6__stcsP6__halfS__nbst = declare_device( + "_ZL6__stcsP6__halfS__nbst", + void(CPointer(CPointer(_type___half)), CPointer(_type___half)), ) - def _ZL12__lows2half27__half2S__1_caller(arg_0, arg_1): - return _ZL12__lows2half27__half2S__1(arg_0, arg_1) + def _ZL6__stcsP6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__stcsP6__halfS__nbst(arg_0, arg_1) - @lower(__lows2half2, _type___half2, _type___half2) + @lower(__stcs, CPointer(_type___half), _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__lows2half27__half2S__1", shim_raw_str - ) + shim_stream.write_with_key("_ZL6__stcsP6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__lows2half27__half2S__1_caller, + _ZL6__stcsP6__halfS__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + void, CPointer(CPointer(_type___half)), CPointer(_type___half) ), ptrs, ) -_lower__ZL12__lows2half27__half2S__1(shim_stream, shim_obj) +_lower__ZL6__stcsP6__halfS__nbst(shim_stream, shim_obj) -def __highs2half2(): +def __stwt(): pass -def _lower__ZL13__highs2half27__half2S__1(shim_stream, shim_obj): +def _lower__ZL6__stwtP6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL13__highs2half27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __highs2half2(*a, *b); + _ZL6__stwtP6__halfS__nbst(int &retval , __half ** ptr, __half* value) { + __stwt(*ptr, *value); return 0; } """ - _ZL13__highs2half27__half2S__1 = declare_device( - "_ZL13__highs2half27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZL6__stwtP6__halfS__nbst = declare_device( + "_ZL6__stwtP6__halfS__nbst", + void(CPointer(CPointer(_type___half)), CPointer(_type___half)), ) - def _ZL13__highs2half27__half2S__1_caller(arg_0, arg_1): - return _ZL13__highs2half27__half2S__1(arg_0, arg_1) + def _ZL6__stwtP6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__stwtP6__halfS__nbst(arg_0, arg_1) - @lower(__highs2half2, _type___half2, _type___half2) + @lower(__stwt, CPointer(_type___half), _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL13__highs2half27__half2S__1", shim_raw_str - ) + shim_stream.write_with_key("_ZL6__stwtP6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL13__highs2half27__half2S__1_caller, + _ZL6__stwtP6__halfS__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + void, CPointer(CPointer(_type___half)), CPointer(_type___half) ), ptrs, ) -_lower__ZL13__highs2half27__half2S__1(shim_stream, shim_obj) +_lower__ZL6__stwtP6__halfS__nbst(shim_stream, shim_obj) -def __high2half(): +def __habs(): pass -def _lower__ZL11__high2half7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__habs6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__high2half7__half2_1(__half &retval , __half2* a) { - retval = __high2half(*a); + _ZL6__habs6__half_nbst(__half &retval , __half* a) { + retval = __habs(*a); return 0; } """ - _ZL11__high2half7__half2_1 = declare_device( - "_ZL11__high2half7__half2_1", _type___half(CPointer(_type___half2)) + _ZL6__habs6__half_nbst = declare_device( + "_ZL6__habs6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL11__high2half7__half2_1_caller(arg_0): - return _ZL11__high2half7__half2_1(arg_0) + def _ZL6__habs6__half_nbst_caller(arg_0): + return _ZL6__habs6__half_nbst(arg_0) - @lower(__high2half, _type___half2) + @lower(__habs, _type___half) + @lower(abs, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__high2half7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__habs6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__high2half7__half2_1_caller, - signature(_type___half, CPointer(_type___half2)), + _ZL6__habs6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL11__high2half7__half2_1(shim_stream, shim_obj) +_lower__ZL6__habs6__half_nbst(shim_stream, shim_obj) -def __low2half(): +def __hadd(): pass -def _lower__ZL10__low2half7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__hadd6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL10__low2half7__half2_1(__half &retval , __half2* a) { - retval = __low2half(*a); + _ZL6__hadd6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hadd(*a, *b); return 0; } """ - _ZL10__low2half7__half2_1 = declare_device( - "_ZL10__low2half7__half2_1", _type___half(CPointer(_type___half2)) + _ZL6__hadd6__halfS__nbst = declare_device( + "_ZL6__hadd6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL10__low2half7__half2_1_caller(arg_0): - return _ZL10__low2half7__half2_1(arg_0) + def _ZL6__hadd6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hadd6__halfS__nbst(arg_0, arg_1) - @lower(__low2half, _type___half2) + @lower(__hadd, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__low2half7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hadd6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL10__low2half7__half2_1_caller, - signature(_type___half, CPointer(_type___half2)), + _ZL6__hadd6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL10__low2half7__half2_1(shim_stream, shim_obj) +_lower__ZL6__hadd6__halfS__nbst(shim_stream, shim_obj) -def __hisinf(): +def __hsub(): pass -def _lower__ZL8__hisinf6__half_1(shim_stream, shim_obj): +def _lower__ZL6__hsub6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL8__hisinf6__half_1(int &retval , __half* a) { - retval = __hisinf(*a); + _ZL6__hsub6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hsub(*a, *b); return 0; } """ - _ZL8__hisinf6__half_1 = declare_device( - "_ZL8__hisinf6__half_1", int32(CPointer(_type___half)) + _ZL6__hsub6__halfS__nbst = declare_device( + "_ZL6__hsub6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL8__hisinf6__half_1_caller(arg_0): - return _ZL8__hisinf6__half_1(arg_0) + def _ZL6__hsub6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hsub6__halfS__nbst(arg_0, arg_1) - @lower(__hisinf, _type___half) + @lower(__hsub, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hisinf6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hsub6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL8__hisinf6__half_1_caller, - signature(int32, CPointer(_type___half)), + _ZL6__hsub6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL8__hisinf6__half_1(shim_stream, shim_obj) +_lower__ZL6__hsub6__halfS__nbst(shim_stream, shim_obj) -def __halves2half2(): +def __hmul(): pass -def _lower__ZL14__halves2half26__halfS__1(shim_stream, shim_obj): +def _lower__ZL6__hmul6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__halves2half26__halfS__1(__half2 &retval , __half* a, __half* b) { - retval = __halves2half2(*a, *b); + _ZL6__hmul6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hmul(*a, *b); return 0; } """ - _ZL14__halves2half26__halfS__1 = declare_device( - "_ZL14__halves2half26__halfS__1", - _type___half2(CPointer(_type___half), CPointer(_type___half)), + _ZL6__hmul6__halfS__nbst = declare_device( + "_ZL6__hmul6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL14__halves2half26__halfS__1_caller(arg_0, arg_1): - return _ZL14__halves2half26__halfS__1(arg_0, arg_1) + def _ZL6__hmul6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hmul6__halfS__nbst(arg_0, arg_1) - @lower(__halves2half2, _type___half, _type___half) + @lower(__hmul, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL14__halves2half26__halfS__1", shim_raw_str - ) + shim_stream.write_with_key("_ZL6__hmul6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__halves2half26__halfS__1_caller, + _ZL6__hmul6__halfS__nbst_caller, signature( - _type___half2, CPointer(_type___half), CPointer(_type___half) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL14__halves2half26__halfS__1(shim_stream, shim_obj) +_lower__ZL6__hmul6__halfS__nbst(shim_stream, shim_obj) -def __low2half2(): +def __hadd_rn(): pass -def _lower__ZL11__low2half27__half2_1(shim_stream, shim_obj): +def _lower__ZL9__hadd_rn6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__low2half27__half2_1(__half2 &retval , __half2* a) { - retval = __low2half2(*a); + _ZL9__hadd_rn6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hadd_rn(*a, *b); return 0; } """ - _ZL11__low2half27__half2_1 = declare_device( - "_ZL11__low2half27__half2_1", _type___half2(CPointer(_type___half2)) + _ZL9__hadd_rn6__halfS__nbst = declare_device( + "_ZL9__hadd_rn6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__low2half27__half2_1_caller(arg_0): - return _ZL11__low2half27__half2_1(arg_0) + def _ZL9__hadd_rn6__halfS__nbst_caller(arg_0, arg_1): + return _ZL9__hadd_rn6__halfS__nbst(arg_0, arg_1) - @lower(__low2half2, _type___half2) + @lower(__hadd_rn, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__low2half27__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL9__hadd_rn6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__low2half27__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), + _ZL9__hadd_rn6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL11__low2half27__half2_1(shim_stream, shim_obj) +_lower__ZL9__hadd_rn6__halfS__nbst(shim_stream, shim_obj) -def __high2half2(): +def __hsub_rn(): pass -def _lower__ZL12__high2half27__half2_1(shim_stream, shim_obj): +def _lower__ZL9__hsub_rn6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL12__high2half27__half2_1(__half2 &retval , __half2* a) { - retval = __high2half2(*a); + _ZL9__hsub_rn6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hsub_rn(*a, *b); return 0; } """ - _ZL12__high2half27__half2_1 = declare_device( - "_ZL12__high2half27__half2_1", _type___half2(CPointer(_type___half2)) + _ZL9__hsub_rn6__halfS__nbst = declare_device( + "_ZL9__hsub_rn6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL12__high2half27__half2_1_caller(arg_0): - return _ZL12__high2half27__half2_1(arg_0) + def _ZL9__hsub_rn6__halfS__nbst_caller(arg_0, arg_1): + return _ZL9__hsub_rn6__halfS__nbst(arg_0, arg_1) - @lower(__high2half2, _type___half2) + @lower(__hsub_rn, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12__high2half27__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL9__hsub_rn6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL12__high2half27__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), + _ZL9__hsub_rn6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL12__high2half27__half2_1(shim_stream, shim_obj) +_lower__ZL9__hsub_rn6__halfS__nbst(shim_stream, shim_obj) -def __half_as_short(): +def __hmul_rn(): pass -def _lower__ZL15__half_as_short6__half_1(shim_stream, shim_obj): +def _lower__ZL9__hmul_rn6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__half_as_short6__half_1(short &retval , __half* h) { - retval = __half_as_short(*h); + _ZL9__hmul_rn6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hmul_rn(*a, *b); return 0; } """ - _ZL15__half_as_short6__half_1 = declare_device( - "_ZL15__half_as_short6__half_1", int16(CPointer(_type___half)) + _ZL9__hmul_rn6__halfS__nbst = declare_device( + "_ZL9__hmul_rn6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL15__half_as_short6__half_1_caller(arg_0): - return _ZL15__half_as_short6__half_1(arg_0) + def _ZL9__hmul_rn6__halfS__nbst_caller(arg_0, arg_1): + return _ZL9__hmul_rn6__halfS__nbst(arg_0, arg_1) - @lower(__half_as_short, _type___half) + @lower(__hmul_rn, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL15__half_as_short6__half_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL9__hmul_rn6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__half_as_short6__half_1_caller, - signature(int16, CPointer(_type___half)), + _ZL9__hmul_rn6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL15__half_as_short6__half_1(shim_stream, shim_obj) +_lower__ZL9__hmul_rn6__halfS__nbst(shim_stream, shim_obj) -def __half_as_ushort(): +def __hdiv(): pass -def _lower__ZL16__half_as_ushort6__half_1(shim_stream, shim_obj): +def _lower__ZL6__hdiv6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__half_as_ushort6__half_1(unsigned short &retval , __half* h) { - retval = __half_as_ushort(*h); + _ZL6__hdiv6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hdiv(*a, *b); return 0; } """ - _ZL16__half_as_ushort6__half_1 = declare_device( - "_ZL16__half_as_ushort6__half_1", uint16(CPointer(_type___half)) + _ZL6__hdiv6__halfS__nbst = declare_device( + "_ZL6__hdiv6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL16__half_as_ushort6__half_1_caller(arg_0): - return _ZL16__half_as_ushort6__half_1(arg_0) + def _ZL6__hdiv6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hdiv6__halfS__nbst(arg_0, arg_1) - @lower(__half_as_ushort, _type___half) + @lower(__hdiv, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL16__half_as_ushort6__half_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL6__hdiv6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__half_as_ushort6__half_1_caller, - signature(uint16, CPointer(_type___half)), + _ZL6__hdiv6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL16__half_as_ushort6__half_1(shim_stream, shim_obj) +_lower__ZL6__hdiv6__halfS__nbst(shim_stream, shim_obj) -def __short_as_half(): +def __hadd_sat(): pass -def _lower__ZL15__short_as_halfs_1(shim_stream, shim_obj): +def _lower__ZL10__hadd_sat6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__short_as_halfs_1(__half &retval , short* i) { - retval = __short_as_half(*i); + _ZL10__hadd_sat6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hadd_sat(*a, *b); return 0; } """ - _ZL15__short_as_halfs_1 = declare_device( - "_ZL15__short_as_halfs_1", _type___half(CPointer(int16)) + _ZL10__hadd_sat6__halfS__nbst = declare_device( + "_ZL10__hadd_sat6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL15__short_as_halfs_1_caller(arg_0): - return _ZL15__short_as_halfs_1(arg_0) + def _ZL10__hadd_sat6__halfS__nbst_caller(arg_0, arg_1): + return _ZL10__hadd_sat6__halfS__nbst(arg_0, arg_1) - @lower(__short_as_half, int16) + @lower(__hadd_sat, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL15__short_as_halfs_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL10__hadd_sat6__halfS__nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__short_as_halfs_1_caller, - signature(_type___half, CPointer(int16)), + _ZL10__hadd_sat6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL15__short_as_halfs_1(shim_stream, shim_obj) +_lower__ZL10__hadd_sat6__halfS__nbst(shim_stream, shim_obj) -def __ushort_as_half(): +def __hsub_sat(): pass -def _lower__ZL16__ushort_as_halft_1(shim_stream, shim_obj): +def _lower__ZL10__hsub_sat6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__ushort_as_halft_1(__half &retval , unsigned short* i) { - retval = __ushort_as_half(*i); + _ZL10__hsub_sat6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hsub_sat(*a, *b); return 0; } """ - _ZL16__ushort_as_halft_1 = declare_device( - "_ZL16__ushort_as_halft_1", _type___half(CPointer(uint16)) + _ZL10__hsub_sat6__halfS__nbst = declare_device( + "_ZL10__hsub_sat6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL16__ushort_as_halft_1_caller(arg_0): - return _ZL16__ushort_as_halft_1(arg_0) + def _ZL10__hsub_sat6__halfS__nbst_caller(arg_0, arg_1): + return _ZL10__hsub_sat6__halfS__nbst(arg_0, arg_1) - @lower(__ushort_as_half, uint16) + @lower(__hsub_sat, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL16__ushort_as_halft_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL10__hsub_sat6__halfS__nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__ushort_as_halft_1_caller, - signature(_type___half, CPointer(uint16)), + _ZL10__hsub_sat6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL16__ushort_as_halft_1(shim_stream, shim_obj) +_lower__ZL10__hsub_sat6__halfS__nbst(shim_stream, shim_obj) -def __hmax(): +def __hmul_sat(): pass -def _lower__ZL6__hmax6__halfS__1(shim_stream, shim_obj): +def _lower__ZL10__hmul_sat6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__hmax6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hmax(*a, *b); + _ZL10__hmul_sat6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hmul_sat(*a, *b); return 0; } """ - _ZL6__hmax6__halfS__1 = declare_device( - "_ZL6__hmax6__halfS__1", + _ZL10__hmul_sat6__halfS__nbst = declare_device( + "_ZL10__hmul_sat6__halfS__nbst", _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__hmax6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hmax6__halfS__1(arg_0, arg_1) + def _ZL10__hmul_sat6__halfS__nbst_caller(arg_0, arg_1): + return _ZL10__hmul_sat6__halfS__nbst(arg_0, arg_1) - @lower(__hmax, _type___half, _type___half) + @lower(__hmul_sat, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hmax6__halfS__1", shim_raw_str) + shim_stream.write_with_key( + "_ZL10__hmul_sat6__halfS__nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__hmax6__halfS__1_caller, + _ZL10__hmul_sat6__halfS__nbst_caller, signature( _type___half, CPointer(_type___half), CPointer(_type___half) ), @@ -4656,8470 +4718,2685 @@ def impl(context, builder, sig, args): ) -_lower__ZL6__hmax6__halfS__1(shim_stream, shim_obj) +_lower__ZL10__hmul_sat6__halfS__nbst(shim_stream, shim_obj) -def __hmin(): +def __hfma(): pass -def _lower__ZL6__hmin6__halfS__1(shim_stream, shim_obj): +def _lower__ZL6__hfma6__halfS_S__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__hmin6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hmin(*a, *b); + _ZL6__hfma6__halfS_S__nbst(__half &retval , __half* a, __half* b, __half* c) { + retval = __hfma(*a, *b, *c); return 0; } """ - _ZL6__hmin6__halfS__1 = declare_device( - "_ZL6__hmin6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), + _ZL6__hfma6__halfS_S__nbst = declare_device( + "_ZL6__hfma6__halfS_S__nbst", + _type___half( + CPointer(_type___half), + CPointer(_type___half), + CPointer(_type___half), + ), ) - def _ZL6__hmin6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hmin6__halfS__1(arg_0, arg_1) + def _ZL6__hfma6__halfS_S__nbst_caller(arg_0, arg_1, arg_2): + return _ZL6__hfma6__halfS_S__nbst(arg_0, arg_1, arg_2) - @lower(__hmin, _type___half, _type___half) + @lower(__hfma, _type___half, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hmin6__halfS__1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hfma6__halfS_S__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__hmin6__halfS__1_caller, + _ZL6__hfma6__halfS_S__nbst_caller, signature( - _type___half, CPointer(_type___half), CPointer(_type___half) + _type___half, + CPointer(_type___half), + CPointer(_type___half), + CPointer(_type___half), ), ptrs, ) -_lower__ZL6__hmin6__halfS__1(shim_stream, shim_obj) +_lower__ZL6__hfma6__halfS_S__nbst(shim_stream, shim_obj) -def __hmax2(): +def __hfma_sat(): pass -def _lower__ZL7__hmax27__half2S__1(shim_stream, shim_obj): +def _lower__ZL10__hfma_sat6__halfS_S__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hmax27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hmax2(*a, *b); + _ZL10__hfma_sat6__halfS_S__nbst(__half &retval , __half* a, __half* b, __half* c) { + retval = __hfma_sat(*a, *b, *c); return 0; } """ - _ZL7__hmax27__half2S__1 = declare_device( - "_ZL7__hmax27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZL10__hfma_sat6__halfS_S__nbst = declare_device( + "_ZL10__hfma_sat6__halfS_S__nbst", + _type___half( + CPointer(_type___half), + CPointer(_type___half), + CPointer(_type___half), + ), ) - def _ZL7__hmax27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hmax27__half2S__1(arg_0, arg_1) + def _ZL10__hfma_sat6__halfS_S__nbst_caller(arg_0, arg_1, arg_2): + return _ZL10__hfma_sat6__halfS_S__nbst(arg_0, arg_1, arg_2) - @lower(__hmax2, _type___half2, _type___half2) + @lower(__hfma_sat, _type___half, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hmax27__half2S__1", shim_raw_str) + shim_stream.write_with_key( + "_ZL10__hfma_sat6__halfS_S__nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hmax27__half2S__1_caller, + _ZL10__hfma_sat6__halfS_S__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, + CPointer(_type___half), + CPointer(_type___half), + CPointer(_type___half), ), ptrs, ) -_lower__ZL7__hmax27__half2S__1(shim_stream, shim_obj) +_lower__ZL10__hfma_sat6__halfS_S__nbst(shim_stream, shim_obj) -def __hmin2(): +def __hneg(): pass -def _lower__ZL7__hmin27__half2S__1(shim_stream, shim_obj): +def _lower__ZL6__hneg6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hmin27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hmin2(*a, *b); + _ZL6__hneg6__half_nbst(__half &retval , __half* a) { + retval = __hneg(*a); return 0; } """ - _ZL7__hmin27__half2S__1 = declare_device( - "_ZL7__hmin27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZL6__hneg6__half_nbst = declare_device( + "_ZL6__hneg6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL7__hmin27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hmin27__half2S__1(arg_0, arg_1) + def _ZL6__hneg6__half_nbst_caller(arg_0): + return _ZL6__hneg6__half_nbst(arg_0) - @lower(__hmin2, _type___half2, _type___half2) + @lower(__hneg, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hmin27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hneg6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hmin27__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), + _ZL6__hneg6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL7__hmin27__half2S__1(shim_stream, shim_obj) +_lower__ZL6__hneg6__half_nbst(shim_stream, shim_obj) -def __shfl_sync(): +def __heq(): pass -def _lower__ZL11__shfl_syncj7__half2ii_1(shim_stream, shim_obj): +def _lower__ZL5__heq6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__shfl_syncj7__half2ii_1(__half2 &retval , unsigned int* mask, __half2* var, int* srcLane, int* width) { - retval = __shfl_sync(*mask, *var, *srcLane, *width); + _ZL5__heq6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __heq(*a, *b); return 0; } """ - _ZL11__shfl_syncj7__half2ii_1 = declare_device( - "_ZL11__shfl_syncj7__half2ii_1", - _type___half2( - CPointer(uint32), - CPointer(_type___half2), - CPointer(int32), - CPointer(int32), - ), + _ZL5__heq6__halfS__nbst = declare_device( + "_ZL5__heq6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__shfl_syncj7__half2ii_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL11__shfl_syncj7__half2ii_1(arg_0, arg_1, arg_2, arg_3) + def _ZL5__heq6__halfS__nbst_caller(arg_0, arg_1): + return _ZL5__heq6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_sync, uint32, _type___half2, int32, int32) + @lower(__heq, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL11__shfl_syncj7__half2ii_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL5__heq6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__shfl_syncj7__half2ii_1_caller, - signature( - _type___half2, - CPointer(uint32), - CPointer(_type___half2), - CPointer(int32), - CPointer(int32), - ), + _ZL5__heq6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__shfl_syncj7__half2ii_1(shim_stream, shim_obj) +_lower__ZL5__heq6__halfS__nbst(shim_stream, shim_obj) -def __shfl_up_sync(): +def __hne(): pass -def _lower__ZL14__shfl_up_syncj7__half2ji_1(shim_stream, shim_obj): +def _lower__ZL5__hne6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__shfl_up_syncj7__half2ji_1(__half2 &retval , unsigned int* mask, __half2* var, unsigned int* delta, int* width) { - retval = __shfl_up_sync(*mask, *var, *delta, *width); + _ZL5__hne6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hne(*a, *b); return 0; } """ - _ZL14__shfl_up_syncj7__half2ji_1 = declare_device( - "_ZL14__shfl_up_syncj7__half2ji_1", - _type___half2( - CPointer(uint32), - CPointer(_type___half2), - CPointer(uint32), - CPointer(int32), - ), + _ZL5__hne6__halfS__nbst = declare_device( + "_ZL5__hne6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL14__shfl_up_syncj7__half2ji_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL14__shfl_up_syncj7__half2ji_1(arg_0, arg_1, arg_2, arg_3) + def _ZL5__hne6__halfS__nbst_caller(arg_0, arg_1): + return _ZL5__hne6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_up_sync, uint32, _type___half2, uint32, int32) + @lower(__hne, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL14__shfl_up_syncj7__half2ji_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL5__hne6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__shfl_up_syncj7__half2ji_1_caller, - signature( - _type___half2, - CPointer(uint32), - CPointer(_type___half2), - CPointer(uint32), - CPointer(int32), - ), + _ZL5__hne6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL14__shfl_up_syncj7__half2ji_1(shim_stream, shim_obj) +_lower__ZL5__hne6__halfS__nbst(shim_stream, shim_obj) -def __shfl_down_sync(): +def __hle(): pass -def _lower__ZL16__shfl_down_syncj7__half2ji_1(shim_stream, shim_obj): +def _lower__ZL5__hle6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__shfl_down_syncj7__half2ji_1(__half2 &retval , unsigned int* mask, __half2* var, unsigned int* delta, int* width) { - retval = __shfl_down_sync(*mask, *var, *delta, *width); + _ZL5__hle6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hle(*a, *b); return 0; } """ - _ZL16__shfl_down_syncj7__half2ji_1 = declare_device( - "_ZL16__shfl_down_syncj7__half2ji_1", - _type___half2( - CPointer(uint32), - CPointer(_type___half2), - CPointer(uint32), - CPointer(int32), - ), + _ZL5__hle6__halfS__nbst = declare_device( + "_ZL5__hle6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL16__shfl_down_syncj7__half2ji_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL16__shfl_down_syncj7__half2ji_1(arg_0, arg_1, arg_2, arg_3) + def _ZL5__hle6__halfS__nbst_caller(arg_0, arg_1): + return _ZL5__hle6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_down_sync, uint32, _type___half2, uint32, int32) + @lower(__hle, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL16__shfl_down_syncj7__half2ji_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL5__hle6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__shfl_down_syncj7__half2ji_1_caller, - signature( - _type___half2, - CPointer(uint32), - CPointer(_type___half2), - CPointer(uint32), - CPointer(int32), - ), + _ZL5__hle6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL16__shfl_down_syncj7__half2ji_1(shim_stream, shim_obj) +_lower__ZL5__hle6__halfS__nbst(shim_stream, shim_obj) -def __shfl_xor_sync(): +def __hge(): pass -def _lower__ZL15__shfl_xor_syncj7__half2ii_1(shim_stream, shim_obj): +def _lower__ZL5__hge6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__shfl_xor_syncj7__half2ii_1(__half2 &retval , unsigned int* mask, __half2* var, int* laneMask, int* width) { - retval = __shfl_xor_sync(*mask, *var, *laneMask, *width); + _ZL5__hge6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hge(*a, *b); return 0; } """ - _ZL15__shfl_xor_syncj7__half2ii_1 = declare_device( - "_ZL15__shfl_xor_syncj7__half2ii_1", - _type___half2( - CPointer(uint32), - CPointer(_type___half2), - CPointer(int32), - CPointer(int32), - ), + _ZL5__hge6__halfS__nbst = declare_device( + "_ZL5__hge6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL15__shfl_xor_syncj7__half2ii_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL15__shfl_xor_syncj7__half2ii_1(arg_0, arg_1, arg_2, arg_3) + def _ZL5__hge6__halfS__nbst_caller(arg_0, arg_1): + return _ZL5__hge6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_xor_sync, uint32, _type___half2, int32, int32) + @lower(__hge, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL15__shfl_xor_syncj7__half2ii_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL5__hge6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__shfl_xor_syncj7__half2ii_1_caller, - signature( - _type___half2, - CPointer(uint32), - CPointer(_type___half2), - CPointer(int32), - CPointer(int32), - ), + _ZL5__hge6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL15__shfl_xor_syncj7__half2ii_1(shim_stream, shim_obj) +_lower__ZL5__hge6__halfS__nbst(shim_stream, shim_obj) + + +def __hlt(): + pass -def _lower__ZL11__shfl_syncj6__halfii_1(shim_stream, shim_obj): +def _lower__ZL5__hlt6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__shfl_syncj6__halfii_1(__half &retval , unsigned int* mask, __half* var, int* srcLane, int* width) { - retval = __shfl_sync(*mask, *var, *srcLane, *width); + _ZL5__hlt6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hlt(*a, *b); return 0; } """ - _ZL11__shfl_syncj6__halfii_1 = declare_device( - "_ZL11__shfl_syncj6__halfii_1", - _type___half( - CPointer(uint32), - CPointer(_type___half), - CPointer(int32), - CPointer(int32), - ), + _ZL5__hlt6__halfS__nbst = declare_device( + "_ZL5__hlt6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__shfl_syncj6__halfii_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL11__shfl_syncj6__halfii_1(arg_0, arg_1, arg_2, arg_3) + def _ZL5__hlt6__halfS__nbst_caller(arg_0, arg_1): + return _ZL5__hlt6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_sync, uint32, _type___half, int32, int32) + @lower(__hlt, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__shfl_syncj6__halfii_1", shim_raw_str) + shim_stream.write_with_key("_ZL5__hlt6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__shfl_syncj6__halfii_1_caller, - signature( - _type___half, - CPointer(uint32), - CPointer(_type___half), - CPointer(int32), - CPointer(int32), - ), + _ZL5__hlt6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__shfl_syncj6__halfii_1(shim_stream, shim_obj) +_lower__ZL5__hlt6__halfS__nbst(shim_stream, shim_obj) + +def __hgt(): + pass -def _lower__ZL14__shfl_up_syncj6__halfji_1(shim_stream, shim_obj): + +def _lower__ZL5__hgt6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL14__shfl_up_syncj6__halfji_1(__half &retval , unsigned int* mask, __half* var, unsigned int* delta, int* width) { - retval = __shfl_up_sync(*mask, *var, *delta, *width); + _ZL5__hgt6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hgt(*a, *b); return 0; } """ - _ZL14__shfl_up_syncj6__halfji_1 = declare_device( - "_ZL14__shfl_up_syncj6__halfji_1", - _type___half( - CPointer(uint32), - CPointer(_type___half), - CPointer(uint32), - CPointer(int32), - ), + _ZL5__hgt6__halfS__nbst = declare_device( + "_ZL5__hgt6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL14__shfl_up_syncj6__halfji_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL14__shfl_up_syncj6__halfji_1(arg_0, arg_1, arg_2, arg_3) + def _ZL5__hgt6__halfS__nbst_caller(arg_0, arg_1): + return _ZL5__hgt6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_up_sync, uint32, _type___half, uint32, int32) + @lower(__hgt, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL14__shfl_up_syncj6__halfji_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL5__hgt6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL14__shfl_up_syncj6__halfji_1_caller, - signature( - _type___half, - CPointer(uint32), - CPointer(_type___half), - CPointer(uint32), - CPointer(int32), - ), + _ZL5__hgt6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL14__shfl_up_syncj6__halfji_1(shim_stream, shim_obj) +_lower__ZL5__hgt6__halfS__nbst(shim_stream, shim_obj) + + +def __hequ(): + pass -def _lower__ZL16__shfl_down_syncj6__halfji_1(shim_stream, shim_obj): +def _lower__ZL6__hequ6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL16__shfl_down_syncj6__halfji_1(__half &retval , unsigned int* mask, __half* var, unsigned int* delta, int* width) { - retval = __shfl_down_sync(*mask, *var, *delta, *width); + _ZL6__hequ6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hequ(*a, *b); return 0; } """ - _ZL16__shfl_down_syncj6__halfji_1 = declare_device( - "_ZL16__shfl_down_syncj6__halfji_1", - _type___half( - CPointer(uint32), - CPointer(_type___half), - CPointer(uint32), - CPointer(int32), - ), + _ZL6__hequ6__halfS__nbst = declare_device( + "_ZL6__hequ6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL16__shfl_down_syncj6__halfji_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL16__shfl_down_syncj6__halfji_1(arg_0, arg_1, arg_2, arg_3) + def _ZL6__hequ6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hequ6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_down_sync, uint32, _type___half, uint32, int32) + @lower(__hequ, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL16__shfl_down_syncj6__halfji_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL6__hequ6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL16__shfl_down_syncj6__halfji_1_caller, - signature( - _type___half, - CPointer(uint32), - CPointer(_type___half), - CPointer(uint32), - CPointer(int32), - ), + _ZL6__hequ6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL16__shfl_down_syncj6__halfji_1(shim_stream, shim_obj) +_lower__ZL6__hequ6__halfS__nbst(shim_stream, shim_obj) + + +def __hneu(): + pass -def _lower__ZL15__shfl_xor_syncj6__halfii_1(shim_stream, shim_obj): +def _lower__ZL6__hneu6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL15__shfl_xor_syncj6__halfii_1(__half &retval , unsigned int* mask, __half* var, int* laneMask, int* width) { - retval = __shfl_xor_sync(*mask, *var, *laneMask, *width); + _ZL6__hneu6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hneu(*a, *b); return 0; } """ - _ZL15__shfl_xor_syncj6__halfii_1 = declare_device( - "_ZL15__shfl_xor_syncj6__halfii_1", - _type___half( - CPointer(uint32), - CPointer(_type___half), - CPointer(int32), - CPointer(int32), - ), + _ZL6__hneu6__halfS__nbst = declare_device( + "_ZL6__hneu6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL15__shfl_xor_syncj6__halfii_1_caller(arg_0, arg_1, arg_2, arg_3): - return _ZL15__shfl_xor_syncj6__halfii_1(arg_0, arg_1, arg_2, arg_3) + def _ZL6__hneu6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hneu6__halfS__nbst(arg_0, arg_1) - @lower(__shfl_xor_sync, uint32, _type___half, int32, int32) + @lower(__hneu, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL15__shfl_xor_syncj6__halfii_1", shim_raw_str - ) + shim_stream.write_with_key("_ZL6__hneu6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL15__shfl_xor_syncj6__halfii_1_caller, - signature( - _type___half, - CPointer(uint32), - CPointer(_type___half), - CPointer(int32), - CPointer(int32), - ), + _ZL6__hneu6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL15__shfl_xor_syncj6__halfii_1(shim_stream, shim_obj) +_lower__ZL6__hneu6__halfS__nbst(shim_stream, shim_obj) -def __ldg(): +def __hleu(): pass -def _lower__ZL5__ldgPK7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__hleu6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL5__ldgPK7__half2_1(__half2 &retval , __half2 ** ptr) { - retval = __ldg(*ptr); + _ZL6__hleu6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hleu(*a, *b); return 0; } """ - _ZL5__ldgPK7__half2_1 = declare_device( - "_ZL5__ldgPK7__half2_1", - _type___half2(CPointer(CPointer(_type___half2))), + _ZL6__hleu6__halfS__nbst = declare_device( + "_ZL6__hleu6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL5__ldgPK7__half2_1_caller(arg_0): - return _ZL5__ldgPK7__half2_1(arg_0) + def _ZL6__hleu6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hleu6__halfS__nbst(arg_0, arg_1) - @lower(__ldg, CPointer(_type___half2)) + @lower(__hleu, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__ldgPK7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hleu6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL5__ldgPK7__half2_1_caller, - signature(_type___half2, CPointer(CPointer(_type___half2))), + _ZL6__hleu6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL5__ldgPK7__half2_1(shim_stream, shim_obj) +_lower__ZL6__hleu6__halfS__nbst(shim_stream, shim_obj) + + +def __hgeu(): + pass -def _lower__ZL5__ldgPK6__half_1(shim_stream, shim_obj): +def _lower__ZL6__hgeu6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL5__ldgPK6__half_1(__half &retval , __half ** ptr) { - retval = __ldg(*ptr); + _ZL6__hgeu6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hgeu(*a, *b); return 0; } """ - _ZL5__ldgPK6__half_1 = declare_device( - "_ZL5__ldgPK6__half_1", _type___half(CPointer(CPointer(_type___half))) + _ZL6__hgeu6__halfS__nbst = declare_device( + "_ZL6__hgeu6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL5__ldgPK6__half_1_caller(arg_0): - return _ZL5__ldgPK6__half_1(arg_0) + def _ZL6__hgeu6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hgeu6__halfS__nbst(arg_0, arg_1) - @lower(__ldg, CPointer(_type___half)) + @lower(__hgeu, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__ldgPK6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hgeu6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL5__ldgPK6__half_1_caller, - signature(_type___half, CPointer(CPointer(_type___half))), + _ZL6__hgeu6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL5__ldgPK6__half_1(shim_stream, shim_obj) +_lower__ZL6__hgeu6__halfS__nbst(shim_stream, shim_obj) -def __ldcg(): +def __hltu(): pass -def _lower__ZL6__ldcgPK7__half2_1(shim_stream, shim_obj): +def _lower__ZL6__hltu6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcgPK7__half2_1(__half2 &retval , __half2 ** ptr) { - retval = __ldcg(*ptr); + _ZL6__hltu6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hltu(*a, *b); return 0; } """ - _ZL6__ldcgPK7__half2_1 = declare_device( - "_ZL6__ldcgPK7__half2_1", - _type___half2(CPointer(CPointer(_type___half2))), + _ZL6__hltu6__halfS__nbst = declare_device( + "_ZL6__hltu6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__ldcgPK7__half2_1_caller(arg_0): - return _ZL6__ldcgPK7__half2_1(arg_0) + def _ZL6__hltu6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hltu6__halfS__nbst(arg_0, arg_1) - @lower(__ldcg, CPointer(_type___half2)) + @lower(__hltu, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcgPK7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hltu6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcgPK7__half2_1_caller, - signature(_type___half2, CPointer(CPointer(_type___half2))), + _ZL6__hltu6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL6__ldcgPK7__half2_1(shim_stream, shim_obj) +_lower__ZL6__hltu6__halfS__nbst(shim_stream, shim_obj) + +def __hgtu(): + pass -def _lower__ZL6__ldcgPK6__half_1(shim_stream, shim_obj): + +def _lower__ZL6__hgtu6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcgPK6__half_1(__half &retval , __half ** ptr) { - retval = __ldcg(*ptr); + _ZL6__hgtu6__halfS__nbst(bool &retval , __half* a, __half* b) { + retval = __hgtu(*a, *b); return 0; } """ - _ZL6__ldcgPK6__half_1 = declare_device( - "_ZL6__ldcgPK6__half_1", _type___half(CPointer(CPointer(_type___half))) + _ZL6__hgtu6__halfS__nbst = declare_device( + "_ZL6__hgtu6__halfS__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__ldcgPK6__half_1_caller(arg_0): - return _ZL6__ldcgPK6__half_1(arg_0) + def _ZL6__hgtu6__halfS__nbst_caller(arg_0, arg_1): + return _ZL6__hgtu6__halfS__nbst(arg_0, arg_1) - @lower(__ldcg, CPointer(_type___half)) + @lower(__hgtu, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcgPK6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL6__hgtu6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcgPK6__half_1_caller, - signature(_type___half, CPointer(CPointer(_type___half))), + _ZL6__hgtu6__halfS__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL6__ldcgPK6__half_1(shim_stream, shim_obj) +_lower__ZL6__hgtu6__halfS__nbst(shim_stream, shim_obj) -def __ldca(): +def __hisnan(): pass -def _lower__ZL6__ldcaPK7__half2_1(shim_stream, shim_obj): +def _lower__ZL8__hisnan6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcaPK7__half2_1(__half2 &retval , __half2 ** ptr) { - retval = __ldca(*ptr); + _ZL8__hisnan6__half_nbst(bool &retval , __half* a) { + retval = __hisnan(*a); return 0; } """ - _ZL6__ldcaPK7__half2_1 = declare_device( - "_ZL6__ldcaPK7__half2_1", - _type___half2(CPointer(CPointer(_type___half2))), + _ZL8__hisnan6__half_nbst = declare_device( + "_ZL8__hisnan6__half_nbst", bool_(CPointer(_type___half)) ) - def _ZL6__ldcaPK7__half2_1_caller(arg_0): - return _ZL6__ldcaPK7__half2_1(arg_0) + def _ZL8__hisnan6__half_nbst_caller(arg_0): + return _ZL8__hisnan6__half_nbst(arg_0) - @lower(__ldca, CPointer(_type___half2)) + @lower(__hisnan, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcaPK7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL8__hisnan6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcaPK7__half2_1_caller, - signature(_type___half2, CPointer(CPointer(_type___half2))), + _ZL8__hisnan6__half_nbst_caller, + signature(bool_, CPointer(_type___half)), ptrs, ) -_lower__ZL6__ldcaPK7__half2_1(shim_stream, shim_obj) +_lower__ZL8__hisnan6__half_nbst(shim_stream, shim_obj) -def _lower__ZL6__ldcaPK6__half_1(shim_stream, shim_obj): +def __hmax_nan(): + pass + + +def _lower__ZL10__hmax_nan6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcaPK6__half_1(__half &retval , __half ** ptr) { - retval = __ldca(*ptr); + _ZL10__hmax_nan6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hmax_nan(*a, *b); return 0; } """ - _ZL6__ldcaPK6__half_1 = declare_device( - "_ZL6__ldcaPK6__half_1", _type___half(CPointer(CPointer(_type___half))) + _ZL10__hmax_nan6__halfS__nbst = declare_device( + "_ZL10__hmax_nan6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__ldcaPK6__half_1_caller(arg_0): - return _ZL6__ldcaPK6__half_1(arg_0) + def _ZL10__hmax_nan6__halfS__nbst_caller(arg_0, arg_1): + return _ZL10__hmax_nan6__halfS__nbst(arg_0, arg_1) - @lower(__ldca, CPointer(_type___half)) + @lower(__hmax_nan, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcaPK6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL10__hmax_nan6__halfS__nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcaPK6__half_1_caller, - signature(_type___half, CPointer(CPointer(_type___half))), + _ZL10__hmax_nan6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL6__ldcaPK6__half_1(shim_stream, shim_obj) +_lower__ZL10__hmax_nan6__halfS__nbst(shim_stream, shim_obj) -def __ldcs(): +def __hmin_nan(): pass -def _lower__ZL6__ldcsPK7__half2_1(shim_stream, shim_obj): +def _lower__ZL10__hmin_nan6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcsPK7__half2_1(__half2 &retval , __half2 ** ptr) { - retval = __ldcs(*ptr); + _ZL10__hmin_nan6__halfS__nbst(__half &retval , __half* a, __half* b) { + retval = __hmin_nan(*a, *b); return 0; } """ - _ZL6__ldcsPK7__half2_1 = declare_device( - "_ZL6__ldcsPK7__half2_1", - _type___half2(CPointer(CPointer(_type___half2))), + _ZL10__hmin_nan6__halfS__nbst = declare_device( + "_ZL10__hmin_nan6__halfS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__ldcsPK7__half2_1_caller(arg_0): - return _ZL6__ldcsPK7__half2_1(arg_0) + def _ZL10__hmin_nan6__halfS__nbst_caller(arg_0, arg_1): + return _ZL10__hmin_nan6__halfS__nbst(arg_0, arg_1) - @lower(__ldcs, CPointer(_type___half2)) + @lower(__hmin_nan, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcsPK7__half2_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL10__hmin_nan6__halfS__nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcsPK7__half2_1_caller, - signature(_type___half2, CPointer(CPointer(_type___half2))), + _ZL10__hmin_nan6__halfS__nbst_caller, + signature( + _type___half, CPointer(_type___half), CPointer(_type___half) + ), ptrs, ) -_lower__ZL6__ldcsPK7__half2_1(shim_stream, shim_obj) +_lower__ZL10__hmin_nan6__halfS__nbst(shim_stream, shim_obj) + + +def __hfma_relu(): + pass -def _lower__ZL6__ldcsPK6__half_1(shim_stream, shim_obj): +def _lower__ZL11__hfma_relu6__halfS_S__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcsPK6__half_1(__half &retval , __half ** ptr) { - retval = __ldcs(*ptr); + _ZL11__hfma_relu6__halfS_S__nbst(__half &retval , __half* a, __half* b, __half* c) { + retval = __hfma_relu(*a, *b, *c); return 0; } """ - _ZL6__ldcsPK6__half_1 = declare_device( - "_ZL6__ldcsPK6__half_1", _type___half(CPointer(CPointer(_type___half))) + _ZL11__hfma_relu6__halfS_S__nbst = declare_device( + "_ZL11__hfma_relu6__halfS_S__nbst", + _type___half( + CPointer(_type___half), + CPointer(_type___half), + CPointer(_type___half), + ), ) - def _ZL6__ldcsPK6__half_1_caller(arg_0): - return _ZL6__ldcsPK6__half_1(arg_0) + def _ZL11__hfma_relu6__halfS_S__nbst_caller(arg_0, arg_1, arg_2): + return _ZL11__hfma_relu6__halfS_S__nbst(arg_0, arg_1, arg_2) - @lower(__ldcs, CPointer(_type___half)) + @lower(__hfma_relu, _type___half, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcsPK6__half_1", shim_raw_str) + shim_stream.write_with_key( + "_ZL11__hfma_relu6__halfS_S__nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcsPK6__half_1_caller, - signature(_type___half, CPointer(CPointer(_type___half))), + _ZL11__hfma_relu6__halfS_S__nbst_caller, + signature( + _type___half, + CPointer(_type___half), + CPointer(_type___half), + CPointer(_type___half), + ), ptrs, ) -_lower__ZL6__ldcsPK6__half_1(shim_stream, shim_obj) +_lower__ZL11__hfma_relu6__halfS_S__nbst(shim_stream, shim_obj) -def __ldlu(): +def hsqrt(): pass -def _lower__ZL6__ldluPK7__half2_1(shim_stream, shim_obj): +def _lower__ZL5hsqrt6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldluPK7__half2_1(__half2 &retval , __half2 ** ptr) { - retval = __ldlu(*ptr); + _ZL5hsqrt6__half_nbst(__half &retval , __half* a) { + retval = hsqrt(*a); return 0; } """ - _ZL6__ldluPK7__half2_1 = declare_device( - "_ZL6__ldluPK7__half2_1", - _type___half2(CPointer(CPointer(_type___half2))), + _ZL5hsqrt6__half_nbst = declare_device( + "_ZL5hsqrt6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__ldluPK7__half2_1_caller(arg_0): - return _ZL6__ldluPK7__half2_1(arg_0) + def _ZL5hsqrt6__half_nbst_caller(arg_0): + return _ZL5hsqrt6__half_nbst(arg_0) - @lower(__ldlu, CPointer(_type___half2)) + @lower(hsqrt, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldluPK7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL5hsqrt6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldluPK7__half2_1_caller, - signature(_type___half2, CPointer(CPointer(_type___half2))), + _ZL5hsqrt6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__ldluPK7__half2_1(shim_stream, shim_obj) +_lower__ZL5hsqrt6__half_nbst(shim_stream, shim_obj) + + +def hrsqrt(): + pass -def _lower__ZL6__ldluPK6__half_1(shim_stream, shim_obj): +def _lower__ZL6hrsqrt6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldluPK6__half_1(__half &retval , __half ** ptr) { - retval = __ldlu(*ptr); + _ZL6hrsqrt6__half_nbst(__half &retval , __half* a) { + retval = hrsqrt(*a); return 0; } """ - _ZL6__ldluPK6__half_1 = declare_device( - "_ZL6__ldluPK6__half_1", _type___half(CPointer(CPointer(_type___half))) + _ZL6hrsqrt6__half_nbst = declare_device( + "_ZL6hrsqrt6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__ldluPK6__half_1_caller(arg_0): - return _ZL6__ldluPK6__half_1(arg_0) + def _ZL6hrsqrt6__half_nbst_caller(arg_0): + return _ZL6hrsqrt6__half_nbst(arg_0) - @lower(__ldlu, CPointer(_type___half)) + @lower(hrsqrt, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldluPK6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL6hrsqrt6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldluPK6__half_1_caller, - signature(_type___half, CPointer(CPointer(_type___half))), + _ZL6hrsqrt6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__ldluPK6__half_1(shim_stream, shim_obj) +_lower__ZL6hrsqrt6__half_nbst(shim_stream, shim_obj) -def __ldcv(): +def hrcp(): pass -def _lower__ZL6__ldcvPK7__half2_1(shim_stream, shim_obj): +def _lower__ZL4hrcp6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcvPK7__half2_1(__half2 &retval , __half2 ** ptr) { - retval = __ldcv(*ptr); + _ZL4hrcp6__half_nbst(__half &retval , __half* a) { + retval = hrcp(*a); return 0; } """ - _ZL6__ldcvPK7__half2_1 = declare_device( - "_ZL6__ldcvPK7__half2_1", - _type___half2(CPointer(CPointer(_type___half2))), + _ZL4hrcp6__half_nbst = declare_device( + "_ZL4hrcp6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__ldcvPK7__half2_1_caller(arg_0): - return _ZL6__ldcvPK7__half2_1(arg_0) + def _ZL4hrcp6__half_nbst_caller(arg_0): + return _ZL4hrcp6__half_nbst(arg_0) - @lower(__ldcv, CPointer(_type___half2)) + @lower(hrcp, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcvPK7__half2_1", shim_raw_str) + shim_stream.write_with_key("_ZL4hrcp6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcvPK7__half2_1_caller, - signature(_type___half2, CPointer(CPointer(_type___half2))), + _ZL4hrcp6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__ldcvPK7__half2_1(shim_stream, shim_obj) +_lower__ZL4hrcp6__half_nbst(shim_stream, shim_obj) + + +def hlog(): + pass -def _lower__ZL6__ldcvPK6__half_1(shim_stream, shim_obj): +def _lower__ZL4hlog6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__ldcvPK6__half_1(__half &retval , __half ** ptr) { - retval = __ldcv(*ptr); + _ZL4hlog6__half_nbst(__half &retval , __half* a) { + retval = hlog(*a); return 0; } """ - _ZL6__ldcvPK6__half_1 = declare_device( - "_ZL6__ldcvPK6__half_1", _type___half(CPointer(CPointer(_type___half))) + _ZL4hlog6__half_nbst = declare_device( + "_ZL4hlog6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__ldcvPK6__half_1_caller(arg_0): - return _ZL6__ldcvPK6__half_1(arg_0) + def _ZL4hlog6__half_nbst_caller(arg_0): + return _ZL4hlog6__half_nbst(arg_0) - @lower(__ldcv, CPointer(_type___half)) + @lower(hlog, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__ldcvPK6__half_1", shim_raw_str) + shim_stream.write_with_key("_ZL4hlog6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__ldcvPK6__half_1_caller, - signature(_type___half, CPointer(CPointer(_type___half))), + _ZL4hlog6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__ldcvPK6__half_1(shim_stream, shim_obj) +_lower__ZL4hlog6__half_nbst(shim_stream, shim_obj) -def __stwb(): +def hlog2(): pass -def _lower__ZL6__stwbP7__half2S__1(shim_stream, shim_obj): +def _lower__ZL5hlog26__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stwbP7__half2S__1(int &retval , __half2 ** ptr, __half2* value) { - __stwb(*ptr, *value); + _ZL5hlog26__half_nbst(__half &retval , __half* a) { + retval = hlog2(*a); return 0; } """ - _ZL6__stwbP7__half2S__1 = declare_device( - "_ZL6__stwbP7__half2S__1", - void(CPointer(CPointer(_type___half2)), CPointer(_type___half2)), + _ZL5hlog26__half_nbst = declare_device( + "_ZL5hlog26__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stwbP7__half2S__1_caller(arg_0, arg_1): - return _ZL6__stwbP7__half2S__1(arg_0, arg_1) + def _ZL5hlog26__half_nbst_caller(arg_0): + return _ZL5hlog26__half_nbst(arg_0) - @lower(__stwb, CPointer(_type___half2), _type___half2) + @lower(hlog2, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stwbP7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZL5hlog26__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stwbP7__half2S__1_caller, - signature( - void, CPointer(CPointer(_type___half2)), CPointer(_type___half2) - ), + _ZL5hlog26__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stwbP7__half2S__1(shim_stream, shim_obj) +_lower__ZL5hlog26__half_nbst(shim_stream, shim_obj) -def _lower__ZL6__stwbP6__halfS__1(shim_stream, shim_obj): +def hlog10(): + pass + + +def _lower__ZL6hlog106__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stwbP6__halfS__1(int &retval , __half ** ptr, __half* value) { - __stwb(*ptr, *value); + _ZL6hlog106__half_nbst(__half &retval , __half* a) { + retval = hlog10(*a); return 0; } """ - _ZL6__stwbP6__halfS__1 = declare_device( - "_ZL6__stwbP6__halfS__1", - void(CPointer(CPointer(_type___half)), CPointer(_type___half)), + _ZL6hlog106__half_nbst = declare_device( + "_ZL6hlog106__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stwbP6__halfS__1_caller(arg_0, arg_1): - return _ZL6__stwbP6__halfS__1(arg_0, arg_1) + def _ZL6hlog106__half_nbst_caller(arg_0): + return _ZL6hlog106__half_nbst(arg_0) - @lower(__stwb, CPointer(_type___half), _type___half) + @lower(hlog10, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stwbP6__halfS__1", shim_raw_str) + shim_stream.write_with_key("_ZL6hlog106__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stwbP6__halfS__1_caller, - signature( - void, CPointer(CPointer(_type___half)), CPointer(_type___half) - ), + _ZL6hlog106__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stwbP6__halfS__1(shim_stream, shim_obj) +_lower__ZL6hlog106__half_nbst(shim_stream, shim_obj) -def __stcg(): +def hexp(): pass -def _lower__ZL6__stcgP7__half2S__1(shim_stream, shim_obj): +def _lower__ZL4hexp6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stcgP7__half2S__1(int &retval , __half2 ** ptr, __half2* value) { - __stcg(*ptr, *value); + _ZL4hexp6__half_nbst(__half &retval , __half* a) { + retval = hexp(*a); return 0; } """ - _ZL6__stcgP7__half2S__1 = declare_device( - "_ZL6__stcgP7__half2S__1", - void(CPointer(CPointer(_type___half2)), CPointer(_type___half2)), + _ZL4hexp6__half_nbst = declare_device( + "_ZL4hexp6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stcgP7__half2S__1_caller(arg_0, arg_1): - return _ZL6__stcgP7__half2S__1(arg_0, arg_1) + def _ZL4hexp6__half_nbst_caller(arg_0): + return _ZL4hexp6__half_nbst(arg_0) - @lower(__stcg, CPointer(_type___half2), _type___half2) + @lower(hexp, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stcgP7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZL4hexp6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stcgP7__half2S__1_caller, - signature( - void, CPointer(CPointer(_type___half2)), CPointer(_type___half2) - ), + _ZL4hexp6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stcgP7__half2S__1(shim_stream, shim_obj) +_lower__ZL4hexp6__half_nbst(shim_stream, shim_obj) + + +def htanh_approx(): + pass -def _lower__ZL6__stcgP6__halfS__1(shim_stream, shim_obj): +def _lower__ZL12htanh_approx6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stcgP6__halfS__1(int &retval , __half ** ptr, __half* value) { - __stcg(*ptr, *value); + _ZL12htanh_approx6__half_nbst(__half &retval , __half* a) { + retval = htanh_approx(*a); return 0; } """ - _ZL6__stcgP6__halfS__1 = declare_device( - "_ZL6__stcgP6__halfS__1", - void(CPointer(CPointer(_type___half)), CPointer(_type___half)), + _ZL12htanh_approx6__half_nbst = declare_device( + "_ZL12htanh_approx6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stcgP6__halfS__1_caller(arg_0, arg_1): - return _ZL6__stcgP6__halfS__1(arg_0, arg_1) + def _ZL12htanh_approx6__half_nbst_caller(arg_0): + return _ZL12htanh_approx6__half_nbst(arg_0) - @lower(__stcg, CPointer(_type___half), _type___half) + @lower(htanh_approx, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stcgP6__halfS__1", shim_raw_str) + shim_stream.write_with_key( + "_ZL12htanh_approx6__half_nbst", shim_raw_str + ) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stcgP6__halfS__1_caller, - signature( - void, CPointer(CPointer(_type___half)), CPointer(_type___half) - ), + _ZL12htanh_approx6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stcgP6__halfS__1(shim_stream, shim_obj) +_lower__ZL12htanh_approx6__half_nbst(shim_stream, shim_obj) -def __stcs(): +def htanh(): pass -def _lower__ZL6__stcsP7__half2S__1(shim_stream, shim_obj): +def _lower__ZL5htanh6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stcsP7__half2S__1(int &retval , __half2 ** ptr, __half2* value) { - __stcs(*ptr, *value); + _ZL5htanh6__half_nbst(__half &retval , __half* a) { + retval = htanh(*a); return 0; } """ - _ZL6__stcsP7__half2S__1 = declare_device( - "_ZL6__stcsP7__half2S__1", - void(CPointer(CPointer(_type___half2)), CPointer(_type___half2)), + _ZL5htanh6__half_nbst = declare_device( + "_ZL5htanh6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stcsP7__half2S__1_caller(arg_0, arg_1): - return _ZL6__stcsP7__half2S__1(arg_0, arg_1) + def _ZL5htanh6__half_nbst_caller(arg_0): + return _ZL5htanh6__half_nbst(arg_0) - @lower(__stcs, CPointer(_type___half2), _type___half2) + @lower(htanh, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stcsP7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZL5htanh6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stcsP7__half2S__1_caller, - signature( - void, CPointer(CPointer(_type___half2)), CPointer(_type___half2) - ), + _ZL5htanh6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stcsP7__half2S__1(shim_stream, shim_obj) +_lower__ZL5htanh6__half_nbst(shim_stream, shim_obj) + + +def hexp2(): + pass -def _lower__ZL6__stcsP6__halfS__1(shim_stream, shim_obj): +def _lower__ZL5hexp26__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stcsP6__halfS__1(int &retval , __half ** ptr, __half* value) { - __stcs(*ptr, *value); + _ZL5hexp26__half_nbst(__half &retval , __half* a) { + retval = hexp2(*a); return 0; } """ - _ZL6__stcsP6__halfS__1 = declare_device( - "_ZL6__stcsP6__halfS__1", - void(CPointer(CPointer(_type___half)), CPointer(_type___half)), + _ZL5hexp26__half_nbst = declare_device( + "_ZL5hexp26__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stcsP6__halfS__1_caller(arg_0, arg_1): - return _ZL6__stcsP6__halfS__1(arg_0, arg_1) + def _ZL5hexp26__half_nbst_caller(arg_0): + return _ZL5hexp26__half_nbst(arg_0) - @lower(__stcs, CPointer(_type___half), _type___half) + @lower(hexp2, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stcsP6__halfS__1", shim_raw_str) + shim_stream.write_with_key("_ZL5hexp26__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stcsP6__halfS__1_caller, - signature( - void, CPointer(CPointer(_type___half)), CPointer(_type___half) - ), + _ZL5hexp26__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stcsP6__halfS__1(shim_stream, shim_obj) +_lower__ZL5hexp26__half_nbst(shim_stream, shim_obj) -def __stwt(): +def hexp10(): pass -def _lower__ZL6__stwtP7__half2S__1(shim_stream, shim_obj): +def _lower__ZL6hexp106__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stwtP7__half2S__1(int &retval , __half2 ** ptr, __half2* value) { - __stwt(*ptr, *value); + _ZL6hexp106__half_nbst(__half &retval , __half* a) { + retval = hexp10(*a); return 0; } """ - _ZL6__stwtP7__half2S__1 = declare_device( - "_ZL6__stwtP7__half2S__1", - void(CPointer(CPointer(_type___half2)), CPointer(_type___half2)), + _ZL6hexp106__half_nbst = declare_device( + "_ZL6hexp106__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stwtP7__half2S__1_caller(arg_0, arg_1): - return _ZL6__stwtP7__half2S__1(arg_0, arg_1) + def _ZL6hexp106__half_nbst_caller(arg_0): + return _ZL6hexp106__half_nbst(arg_0) - @lower(__stwt, CPointer(_type___half2), _type___half2) + @lower(hexp10, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stwtP7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZL6hexp106__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stwtP7__half2S__1_caller, - signature( - void, CPointer(CPointer(_type___half2)), CPointer(_type___half2) - ), + _ZL6hexp106__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stwtP7__half2S__1(shim_stream, shim_obj) +_lower__ZL6hexp106__half_nbst(shim_stream, shim_obj) + + +def hcos(): + pass -def _lower__ZL6__stwtP6__halfS__1(shim_stream, shim_obj): +def _lower__ZL4hcos6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__stwtP6__halfS__1(int &retval , __half ** ptr, __half* value) { - __stwt(*ptr, *value); + _ZL4hcos6__half_nbst(__half &retval , __half* a) { + retval = hcos(*a); return 0; } """ - _ZL6__stwtP6__halfS__1 = declare_device( - "_ZL6__stwtP6__halfS__1", - void(CPointer(CPointer(_type___half)), CPointer(_type___half)), + _ZL4hcos6__half_nbst = declare_device( + "_ZL4hcos6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__stwtP6__halfS__1_caller(arg_0, arg_1): - return _ZL6__stwtP6__halfS__1(arg_0, arg_1) + def _ZL4hcos6__half_nbst_caller(arg_0): + return _ZL4hcos6__half_nbst(arg_0) - @lower(__stwt, CPointer(_type___half), _type___half) + @lower(hcos, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__stwtP6__halfS__1", shim_raw_str) + shim_stream.write_with_key("_ZL4hcos6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__stwtP6__halfS__1_caller, - signature( - void, CPointer(CPointer(_type___half)), CPointer(_type___half) - ), + _ZL4hcos6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__stwtP6__halfS__1(shim_stream, shim_obj) +_lower__ZL4hcos6__half_nbst(shim_stream, shim_obj) -def __heq2(): +def hsin(): pass -def _lower__ZL6__heq27__half2S__1(shim_stream, shim_obj): +def _lower__ZL4hsin6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__heq27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __heq2(*a, *b); + _ZL4hsin6__half_nbst(__half &retval , __half* a) { + retval = hsin(*a); return 0; } """ - _ZL6__heq27__half2S__1 = declare_device( - "_ZL6__heq27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZL4hsin6__half_nbst = declare_device( + "_ZL4hsin6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL6__heq27__half2S__1_caller(arg_0, arg_1): - return _ZL6__heq27__half2S__1(arg_0, arg_1) + def _ZL4hsin6__half_nbst_caller(arg_0): + return _ZL4hsin6__half_nbst(arg_0) - @lower(__heq2, _type___half2, _type___half2) + @lower(hsin, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__heq27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZL4hsin6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__heq27__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), + _ZL4hsin6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL6__heq27__half2S__1(shim_stream, shim_obj) +_lower__ZL4hsin6__half_nbst(shim_stream, shim_obj) -def __hne2(): +def atomicAdd(): pass -def _lower__ZL6__hne27__half2S__1(shim_stream, shim_obj): +def _lower__ZL9atomicAddP6__halfS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__hne27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hne2(*a, *b); + _ZL9atomicAddP6__halfS__nbst(__half &retval , __half ** address, __half* val) { + retval = atomicAdd(*address, *val); return 0; } """ - _ZL6__hne27__half2S__1 = declare_device( - "_ZL6__hne27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZL9atomicAddP6__halfS__nbst = declare_device( + "_ZL9atomicAddP6__halfS__nbst", + _type___half(CPointer(CPointer(_type___half)), CPointer(_type___half)), ) - def _ZL6__hne27__half2S__1_caller(arg_0, arg_1): - return _ZL6__hne27__half2S__1(arg_0, arg_1) + def _ZL9atomicAddP6__halfS__nbst_caller(arg_0, arg_1): + return _ZL9atomicAddP6__halfS__nbst(arg_0, arg_1) - @lower(__hne2, _type___half2, _type___half2) + @lower(atomicAdd, CPointer(_type___half), _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hne27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZL9atomicAddP6__halfS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__hne27__half2S__1_caller, + _ZL9atomicAddP6__halfS__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, + CPointer(CPointer(_type___half)), + CPointer(_type___half), ), ptrs, ) -_lower__ZL6__hne27__half2S__1(shim_stream, shim_obj) - - -def __hle2(): - pass +_lower__ZL9atomicAddP6__halfS__nbst(shim_stream, shim_obj) -def _lower__ZL6__hle27__half2S__1(shim_stream, shim_obj): +def _lower__ZplRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__hle27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hle2(*a, *b); + _ZplRK6__halfS1__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator+(*lh, *rh); return 0; } """ - _ZL6__hle27__half2S__1 = declare_device( - "_ZL6__hle27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZplRK6__halfS1__nbst = declare_device( + "_ZplRK6__halfS1__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__hle27__half2S__1_caller(arg_0, arg_1): - return _ZL6__hle27__half2S__1(arg_0, arg_1) + def _ZplRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZplRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hle2, _type___half2, _type___half2) + @lower(operator.add, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hle27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZplRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__hle27__half2S__1_caller, + _ZplRK6__halfS1__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL6__hle27__half2S__1(shim_stream, shim_obj) - - -def __hge2(): - pass +_lower__ZplRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL6__hge27__half2S__1(shim_stream, shim_obj): +def _lower__ZmiRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__hge27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hge2(*a, *b); + _ZmiRK6__halfS1__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator-(*lh, *rh); return 0; } """ - _ZL6__hge27__half2S__1 = declare_device( - "_ZL6__hge27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZmiRK6__halfS1__nbst = declare_device( + "_ZmiRK6__halfS1__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__hge27__half2S__1_caller(arg_0, arg_1): - return _ZL6__hge27__half2S__1(arg_0, arg_1) + def _ZmiRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZmiRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hge2, _type___half2, _type___half2) + @lower(operator.sub, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hge27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZmiRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__hge27__half2S__1_caller, + _ZmiRK6__halfS1__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL6__hge27__half2S__1(shim_stream, shim_obj) - - -def __hlt2(): - pass +_lower__ZmiRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL6__hlt27__half2S__1(shim_stream, shim_obj): +def _lower__ZmlRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__hlt27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hlt2(*a, *b); + _ZmlRK6__halfS1__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator*(*lh, *rh); return 0; } """ - _ZL6__hlt27__half2S__1 = declare_device( - "_ZL6__hlt27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZmlRK6__halfS1__nbst = declare_device( + "_ZmlRK6__halfS1__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__hlt27__half2S__1_caller(arg_0, arg_1): - return _ZL6__hlt27__half2S__1(arg_0, arg_1) + def _ZmlRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZmlRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hlt2, _type___half2, _type___half2) + @lower(operator.mul, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hlt27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZmlRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__hlt27__half2S__1_caller, + _ZmlRK6__halfS1__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL6__hlt27__half2S__1(shim_stream, shim_obj) +_lower__ZmlRK6__halfS1__nbst(shim_stream, shim_obj) -def __hgt2(): - pass - - -def _lower__ZL6__hgt27__half2S__1(shim_stream, shim_obj): +def _lower__ZdvRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL6__hgt27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hgt2(*a, *b); + _ZdvRK6__halfS1__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator/(*lh, *rh); return 0; } """ - _ZL6__hgt27__half2S__1 = declare_device( - "_ZL6__hgt27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZdvRK6__halfS1__nbst = declare_device( + "_ZdvRK6__halfS1__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL6__hgt27__half2S__1_caller(arg_0, arg_1): - return _ZL6__hgt27__half2S__1(arg_0, arg_1) + def _ZdvRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZdvRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hgt2, _type___half2, _type___half2) + @lower(operator.truediv, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hgt27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZdvRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL6__hgt27__half2S__1_caller, + _ZdvRK6__halfS1__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL6__hgt27__half2S__1(shim_stream, shim_obj) - - -def __hequ2(): - pass +_lower__ZdvRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL7__hequ27__half2S__1(shim_stream, shim_obj): +def _lower__ZpLR6__halfRKS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hequ27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hequ2(*a, *b); + _ZpLR6__halfRKS__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator+=(*lh, *rh); return 0; } """ - _ZL7__hequ27__half2S__1 = declare_device( - "_ZL7__hequ27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZpLR6__halfRKS__nbst = declare_device( + "_ZpLR6__halfRKS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL7__hequ27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hequ27__half2S__1(arg_0, arg_1) + def _ZpLR6__halfRKS__nbst_caller(arg_0, arg_1): + return _ZpLR6__halfRKS__nbst(arg_0, arg_1) - @lower(__hequ2, _type___half2, _type___half2) + @lower(operator.iadd, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hequ27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZpLR6__halfRKS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hequ27__half2S__1_caller, + _ZpLR6__halfRKS__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL7__hequ27__half2S__1(shim_stream, shim_obj) - - -def __hneu2(): - pass +_lower__ZpLR6__halfRKS__nbst(shim_stream, shim_obj) -def _lower__ZL7__hneu27__half2S__1(shim_stream, shim_obj): +def _lower__ZmIR6__halfRKS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hneu27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hneu2(*a, *b); + _ZmIR6__halfRKS__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator-=(*lh, *rh); return 0; } """ - _ZL7__hneu27__half2S__1 = declare_device( - "_ZL7__hneu27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZmIR6__halfRKS__nbst = declare_device( + "_ZmIR6__halfRKS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL7__hneu27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hneu27__half2S__1(arg_0, arg_1) + def _ZmIR6__halfRKS__nbst_caller(arg_0, arg_1): + return _ZmIR6__halfRKS__nbst(arg_0, arg_1) - @lower(__hneu2, _type___half2, _type___half2) + @lower(operator.isub, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hneu27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZmIR6__halfRKS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hneu27__half2S__1_caller, + _ZmIR6__halfRKS__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL7__hneu27__half2S__1(shim_stream, shim_obj) - - -def __hleu2(): - pass +_lower__ZmIR6__halfRKS__nbst(shim_stream, shim_obj) -def _lower__ZL7__hleu27__half2S__1(shim_stream, shim_obj): +def _lower__ZmLR6__halfRKS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hleu27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hleu2(*a, *b); + _ZmLR6__halfRKS__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator*=(*lh, *rh); return 0; } """ - _ZL7__hleu27__half2S__1 = declare_device( - "_ZL7__hleu27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZmLR6__halfRKS__nbst = declare_device( + "_ZmLR6__halfRKS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL7__hleu27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hleu27__half2S__1(arg_0, arg_1) + def _ZmLR6__halfRKS__nbst_caller(arg_0, arg_1): + return _ZmLR6__halfRKS__nbst(arg_0, arg_1) - @lower(__hleu2, _type___half2, _type___half2) + @lower(operator.imul, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hleu27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZmLR6__halfRKS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hleu27__half2S__1_caller, + _ZmLR6__halfRKS__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL7__hleu27__half2S__1(shim_stream, shim_obj) +_lower__ZmLR6__halfRKS__nbst(shim_stream, shim_obj) -def __hgeu2(): - pass - - -def _lower__ZL7__hgeu27__half2S__1(shim_stream, shim_obj): +def _lower__ZdVR6__halfRKS__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hgeu27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hgeu2(*a, *b); + _ZdVR6__halfRKS__nbst(__half &retval , __half* lh, __half* rh) { + retval = operator/=(*lh, *rh); return 0; } """ - _ZL7__hgeu27__half2S__1 = declare_device( - "_ZL7__hgeu27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZdVR6__halfRKS__nbst = declare_device( + "_ZdVR6__halfRKS__nbst", + _type___half(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL7__hgeu27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hgeu27__half2S__1(arg_0, arg_1) + def _ZdVR6__halfRKS__nbst_caller(arg_0, arg_1): + return _ZdVR6__halfRKS__nbst(arg_0, arg_1) - @lower(__hgeu2, _type___half2, _type___half2) + @lower(operator.itruediv, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hgeu27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZdVR6__halfRKS__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hgeu27__half2S__1_caller, + _ZdVR6__halfRKS__nbst_caller, signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) + _type___half, CPointer(_type___half), CPointer(_type___half) ), ptrs, ) -_lower__ZL7__hgeu27__half2S__1(shim_stream, shim_obj) - - -def __hltu2(): - pass +_lower__ZdVR6__halfRKS__nbst(shim_stream, shim_obj) -def _lower__ZL7__hltu27__half2S__1(shim_stream, shim_obj): +def _lower__ZpsRK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hltu27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hltu2(*a, *b); + _ZpsRK6__half_nbst(__half &retval , __half* h) { + retval = operator+(*h); return 0; } """ - _ZL7__hltu27__half2S__1 = declare_device( - "_ZL7__hltu27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZpsRK6__half_nbst = declare_device( + "_ZpsRK6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL7__hltu27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hltu27__half2S__1(arg_0, arg_1) + def _ZpsRK6__half_nbst_caller(arg_0): + return _ZpsRK6__half_nbst(arg_0) - @lower(__hltu2, _type___half2, _type___half2) + @lower(operator.pos, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hltu27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZpsRK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hltu27__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), + _ZpsRK6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL7__hltu27__half2S__1(shim_stream, shim_obj) - - -def __hgtu2(): - pass +_lower__ZpsRK6__half_nbst(shim_stream, shim_obj) -def _lower__ZL7__hgtu27__half2S__1(shim_stream, shim_obj): +def _lower__ZngRK6__half_nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL7__hgtu27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hgtu2(*a, *b); + _ZngRK6__half_nbst(__half &retval , __half* h) { + retval = operator-(*h); return 0; } """ - _ZL7__hgtu27__half2S__1 = declare_device( - "_ZL7__hgtu27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), + _ZngRK6__half_nbst = declare_device( + "_ZngRK6__half_nbst", _type___half(CPointer(_type___half)) ) - def _ZL7__hgtu27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hgtu27__half2S__1(arg_0, arg_1) + def _ZngRK6__half_nbst_caller(arg_0): + return _ZngRK6__half_nbst(arg_0) - @lower(__hgtu2, _type___half2, _type___half2) + @lower(operator.neg, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hgtu27__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZngRK6__half_nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL7__hgtu27__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), + _ZngRK6__half_nbst_caller, + signature(_type___half, CPointer(_type___half)), ptrs, ) -_lower__ZL7__hgtu27__half2S__1(shim_stream, shim_obj) - - -def __heq2_mask(): - pass +_lower__ZngRK6__half_nbst(shim_stream, shim_obj) -def _lower__ZL11__heq2_mask7__half2S__1(shim_stream, shim_obj): +def _lower__ZeqRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__heq2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __heq2_mask(*a, *b); + _ZeqRK6__halfS1__nbst(bool &retval , __half* lh, __half* rh) { + retval = operator==(*lh, *rh); return 0; } """ - _ZL11__heq2_mask7__half2S__1 = declare_device( - "_ZL11__heq2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), + _ZeqRK6__halfS1__nbst = declare_device( + "_ZeqRK6__halfS1__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__heq2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL11__heq2_mask7__half2S__1(arg_0, arg_1) + def _ZeqRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZeqRK6__halfS1__nbst(arg_0, arg_1) - @lower(__heq2_mask, _type___half2, _type___half2) + @lower(operator.eq, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__heq2_mask7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZeqRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__heq2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), + _ZeqRK6__halfS1__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__heq2_mask7__half2S__1(shim_stream, shim_obj) - - -def __hne2_mask(): - pass +_lower__ZeqRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL11__hne2_mask7__half2S__1(shim_stream, shim_obj): +def _lower__ZneRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__hne2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hne2_mask(*a, *b); + _ZneRK6__halfS1__nbst(bool &retval , __half* lh, __half* rh) { + retval = operator!=(*lh, *rh); return 0; } """ - _ZL11__hne2_mask7__half2S__1 = declare_device( - "_ZL11__hne2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), + _ZneRK6__halfS1__nbst = declare_device( + "_ZneRK6__halfS1__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__hne2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hne2_mask7__half2S__1(arg_0, arg_1) + def _ZneRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZneRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hne2_mask, _type___half2, _type___half2) + @lower(operator.ne, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hne2_mask7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZneRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__hne2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), + _ZneRK6__halfS1__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__hne2_mask7__half2S__1(shim_stream, shim_obj) - - -def __hle2_mask(): - pass +_lower__ZneRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL11__hle2_mask7__half2S__1(shim_stream, shim_obj): +def _lower__ZgtRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__hle2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hle2_mask(*a, *b); + _ZgtRK6__halfS1__nbst(bool &retval , __half* lh, __half* rh) { + retval = operator>(*lh, *rh); return 0; } """ - _ZL11__hle2_mask7__half2S__1 = declare_device( - "_ZL11__hle2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), + _ZgtRK6__halfS1__nbst = declare_device( + "_ZgtRK6__halfS1__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__hle2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hle2_mask7__half2S__1(arg_0, arg_1) + def _ZgtRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZgtRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hle2_mask, _type___half2, _type___half2) + @lower(operator.gt, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hle2_mask7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZgtRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__hle2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), + _ZgtRK6__halfS1__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__hle2_mask7__half2S__1(shim_stream, shim_obj) +_lower__ZgtRK6__halfS1__nbst(shim_stream, shim_obj) -def __hge2_mask(): - pass - - -def _lower__ZL11__hge2_mask7__half2S__1(shim_stream, shim_obj): +def _lower__ZltRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__hge2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hge2_mask(*a, *b); + _ZltRK6__halfS1__nbst(bool &retval , __half* lh, __half* rh) { + retval = operator<(*lh, *rh); return 0; } """ - _ZL11__hge2_mask7__half2S__1 = declare_device( - "_ZL11__hge2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), + _ZltRK6__halfS1__nbst = declare_device( + "_ZltRK6__halfS1__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__hge2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hge2_mask7__half2S__1(arg_0, arg_1) + def _ZltRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZltRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hge2_mask, _type___half2, _type___half2) + @lower(operator.lt, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hge2_mask7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZltRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__hge2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), + _ZltRK6__halfS1__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__hge2_mask7__half2S__1(shim_stream, shim_obj) - - -def __hlt2_mask(): - pass +_lower__ZltRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL11__hlt2_mask7__half2S__1(shim_stream, shim_obj): +def _lower__ZgeRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__hlt2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hlt2_mask(*a, *b); + _ZgeRK6__halfS1__nbst(bool &retval , __half* lh, __half* rh) { + retval = operator>=(*lh, *rh); return 0; } """ - _ZL11__hlt2_mask7__half2S__1 = declare_device( - "_ZL11__hlt2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), + _ZgeRK6__halfS1__nbst = declare_device( + "_ZgeRK6__halfS1__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__hlt2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hlt2_mask7__half2S__1(arg_0, arg_1) + def _ZgeRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZgeRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hlt2_mask, _type___half2, _type___half2) + @lower(operator.ge, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hlt2_mask7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZgeRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__hlt2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), + _ZgeRK6__halfS1__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__hlt2_mask7__half2S__1(shim_stream, shim_obj) - - -def __hgt2_mask(): - pass +_lower__ZgeRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL11__hgt2_mask7__half2S__1(shim_stream, shim_obj): +def _lower__ZleRK6__halfS1__nbst(shim_stream, shim_obj): shim_raw_str = """ extern "C" __device__ int - _ZL11__hgt2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hgt2_mask(*a, *b); + _ZleRK6__halfS1__nbst(bool &retval , __half* lh, __half* rh) { + retval = operator<=(*lh, *rh); return 0; } """ - _ZL11__hgt2_mask7__half2S__1 = declare_device( - "_ZL11__hgt2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), + _ZleRK6__halfS1__nbst = declare_device( + "_ZleRK6__halfS1__nbst", + bool_(CPointer(_type___half), CPointer(_type___half)), ) - def _ZL11__hgt2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hgt2_mask7__half2S__1(arg_0, arg_1) + def _ZleRK6__halfS1__nbst_caller(arg_0, arg_1): + return _ZleRK6__halfS1__nbst(arg_0, arg_1) - @lower(__hgt2_mask, _type___half2, _type___half2) + @lower(operator.le, _type___half, _type___half) def impl(context, builder, sig, args): context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hgt2_mask7__half2S__1", shim_raw_str) + shim_stream.write_with_key("_ZleRK6__halfS1__nbst", shim_raw_str) ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] for ptr, ty, arg in zip(ptrs, sig.args, args): builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) return context.compile_internal( builder, - _ZL11__hgt2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), + _ZleRK6__halfS1__nbst_caller, + signature(bool_, CPointer(_type___half), CPointer(_type___half)), ptrs, ) -_lower__ZL11__hgt2_mask7__half2S__1(shim_stream, shim_obj) - - -def __hequ2_mask(): - pass +_lower__ZleRK6__halfS1__nbst(shim_stream, shim_obj) -def _lower__ZL12__hequ2_mask7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12__hequ2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hequ2_mask(*a, *b); - return 0; - } - """ +@register +class _typing___double2half(ConcreteTemplate): + key = globals()["__double2half"] + cases = [signature(_type___half, float64)] - _ZL12__hequ2_mask7__half2S__1 = declare_device( - "_ZL12__hequ2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), - ) - def _ZL12__hequ2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL12__hequ2_mask7__half2S__1(arg_0, arg_1) +register_global(__double2half, types.Function(_typing___double2half)) - @lower(__hequ2_mask, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__hequ2_mask7__half2S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - return context.compile_internal( - builder, - _ZL12__hequ2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) +@register +class _typing___float2half(ConcreteTemplate): + key = globals()["__float2half"] + cases = [signature(_type___half, float32)] -_lower__ZL12__hequ2_mask7__half2S__1(shim_stream, shim_obj) +register_global(__float2half, types.Function(_typing___float2half)) -def __hneu2_mask(): - pass +@register +class _typing___float2half_rn(ConcreteTemplate): + key = globals()["__float2half_rn"] + cases = [signature(_type___half, float32)] -def _lower__ZL12__hneu2_mask7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12__hneu2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hneu2_mask(*a, *b); - return 0; - } - """ +register_global(__float2half_rn, types.Function(_typing___float2half_rn)) - _ZL12__hneu2_mask7__half2S__1 = declare_device( - "_ZL12__hneu2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), - ) - def _ZL12__hneu2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL12__hneu2_mask7__half2S__1(arg_0, arg_1) +@register +class _typing___float2half_rz(ConcreteTemplate): + key = globals()["__float2half_rz"] + cases = [signature(_type___half, float32)] - @lower(__hneu2_mask, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__hneu2_mask7__half2S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - return context.compile_internal( - builder, - _ZL12__hneu2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) +register_global(__float2half_rz, types.Function(_typing___float2half_rz)) -_lower__ZL12__hneu2_mask7__half2S__1(shim_stream, shim_obj) +@register +class _typing___float2half_rd(ConcreteTemplate): + key = globals()["__float2half_rd"] + cases = [signature(_type___half, float32)] -def __hleu2_mask(): - pass +register_global(__float2half_rd, types.Function(_typing___float2half_rd)) -def _lower__ZL12__hleu2_mask7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12__hleu2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hleu2_mask(*a, *b); - return 0; - } - """ +@register +class _typing___float2half_ru(ConcreteTemplate): + key = globals()["__float2half_ru"] + cases = [signature(_type___half, float32)] - _ZL12__hleu2_mask7__half2S__1 = declare_device( - "_ZL12__hleu2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), - ) - def _ZL12__hleu2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL12__hleu2_mask7__half2S__1(arg_0, arg_1) +register_global(__float2half_ru, types.Function(_typing___float2half_ru)) - @lower(__hleu2_mask, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__hleu2_mask7__half2S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - return context.compile_internal( - builder, - _ZL12__hleu2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) +@register +class _typing___half2float(ConcreteTemplate): + key = globals()["__half2float"] + cases = [signature(float32, _type___half)] -_lower__ZL12__hleu2_mask7__half2S__1(shim_stream, shim_obj) +register_global(__half2float, types.Function(_typing___half2float)) -def __hgeu2_mask(): - pass +@register +class _typing___half2char_rz(ConcreteTemplate): + key = globals()["__half2char_rz"] + cases = [signature(int8, _type___half)] -def _lower__ZL12__hgeu2_mask7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12__hgeu2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hgeu2_mask(*a, *b); - return 0; - } - """ +register_global(__half2char_rz, types.Function(_typing___half2char_rz)) - _ZL12__hgeu2_mask7__half2S__1 = declare_device( - "_ZL12__hgeu2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), - ) - def _ZL12__hgeu2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL12__hgeu2_mask7__half2S__1(arg_0, arg_1) +@register +class _typing___half2uchar_rz(ConcreteTemplate): + key = globals()["__half2uchar_rz"] + cases = [signature(uint8, _type___half)] - @lower(__hgeu2_mask, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__hgeu2_mask7__half2S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - return context.compile_internal( - builder, - _ZL12__hgeu2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) +register_global(__half2uchar_rz, types.Function(_typing___half2uchar_rz)) -_lower__ZL12__hgeu2_mask7__half2S__1(shim_stream, shim_obj) +@register +class _typing___half2short_rz(ConcreteTemplate): + key = globals()["__half2short_rz"] + cases = [signature(int16, _type___half)] -def __hltu2_mask(): - pass +register_global(__half2short_rz, types.Function(_typing___half2short_rz)) -def _lower__ZL12__hltu2_mask7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12__hltu2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hltu2_mask(*a, *b); - return 0; - } - """ +@register +class _typing___half2ushort_rz(ConcreteTemplate): + key = globals()["__half2ushort_rz"] + cases = [signature(uint16, _type___half)] - _ZL12__hltu2_mask7__half2S__1 = declare_device( - "_ZL12__hltu2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), - ) - def _ZL12__hltu2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL12__hltu2_mask7__half2S__1(arg_0, arg_1) +register_global(__half2ushort_rz, types.Function(_typing___half2ushort_rz)) - @lower(__hltu2_mask, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__hltu2_mask7__half2S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - return context.compile_internal( - builder, - _ZL12__hltu2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) +@register +class _typing___half2int_rz(ConcreteTemplate): + key = globals()["__half2int_rz"] + cases = [signature(int32, _type___half)] -_lower__ZL12__hltu2_mask7__half2S__1(shim_stream, shim_obj) +register_global(__half2int_rz, types.Function(_typing___half2int_rz)) -def __hgtu2_mask(): - pass +@register +class _typing___half2uint_rz(ConcreteTemplate): + key = globals()["__half2uint_rz"] + cases = [signature(uint32, _type___half)] -def _lower__ZL12__hgtu2_mask7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12__hgtu2_mask7__half2S__1(unsigned int &retval , __half2* a, __half2* b) { - retval = __hgtu2_mask(*a, *b); - return 0; - } - """ +register_global(__half2uint_rz, types.Function(_typing___half2uint_rz)) - _ZL12__hgtu2_mask7__half2S__1 = declare_device( - "_ZL12__hgtu2_mask7__half2S__1", - uint32(CPointer(_type___half2), CPointer(_type___half2)), - ) - def _ZL12__hgtu2_mask7__half2S__1_caller(arg_0, arg_1): - return _ZL12__hgtu2_mask7__half2S__1(arg_0, arg_1) +@register +class _typing___half2ll_rz(ConcreteTemplate): + key = globals()["__half2ll_rz"] + cases = [signature(int64, _type___half)] - @lower(__hgtu2_mask, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__hgtu2_mask7__half2S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - return context.compile_internal( - builder, - _ZL12__hgtu2_mask7__half2S__1_caller, - signature(uint32, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) +register_global(__half2ll_rz, types.Function(_typing___half2ll_rz)) -_lower__ZL12__hgtu2_mask7__half2S__1(shim_stream, shim_obj) +@register +class _typing___half2ull_rz(ConcreteTemplate): + key = globals()["__half2ull_rz"] + cases = [signature(uint64, _type___half)] -def __hisnan2(): - pass +register_global(__half2ull_rz, types.Function(_typing___half2ull_rz)) -def _lower__ZL9__hisnan27__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL9__hisnan27__half2_1(__half2 &retval , __half2* a) { - retval = __hisnan2(*a); - return 0; - } - """ +@register +class _typing___half2int_rn(ConcreteTemplate): + key = globals()["__half2int_rn"] + cases = [signature(int32, _type___half)] - _ZL9__hisnan27__half2_1 = declare_device( - "_ZL9__hisnan27__half2_1", _type___half2(CPointer(_type___half2)) - ) - def _ZL9__hisnan27__half2_1_caller(arg_0): - return _ZL9__hisnan27__half2_1(arg_0) +register_global(__half2int_rn, types.Function(_typing___half2int_rn)) - @lower(__hisnan2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL9__hisnan27__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - return context.compile_internal( - builder, - _ZL9__hisnan27__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) +@register +class _typing___half2int_rd(ConcreteTemplate): + key = globals()["__half2int_rd"] + cases = [signature(int32, _type___half)] -_lower__ZL9__hisnan27__half2_1(shim_stream, shim_obj) - - -def __hadd2(): - pass - - -def _lower__ZL7__hadd27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hadd27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hadd2(*a, *b); - return 0; - } - """ - - _ZL7__hadd27__half2S__1 = declare_device( - "_ZL7__hadd27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hadd27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hadd27__half2S__1(arg_0, arg_1) - - @lower(__hadd2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hadd27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hadd27__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL7__hadd27__half2S__1(shim_stream, shim_obj) - - -def __hsub2(): - pass - - -def _lower__ZL7__hsub27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hsub27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hsub2(*a, *b); - return 0; - } - """ - - _ZL7__hsub27__half2S__1 = declare_device( - "_ZL7__hsub27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hsub27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hsub27__half2S__1(arg_0, arg_1) - - @lower(__hsub2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hsub27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hsub27__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL7__hsub27__half2S__1(shim_stream, shim_obj) - - -def __hmul2(): - pass - - -def _lower__ZL7__hmul27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hmul27__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hmul2(*a, *b); - return 0; - } - """ - - _ZL7__hmul27__half2S__1 = declare_device( - "_ZL7__hmul27__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hmul27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hmul27__half2S__1(arg_0, arg_1) - - @lower(__hmul2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hmul27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hmul27__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL7__hmul27__half2S__1(shim_stream, shim_obj) - - -def __hadd2_rn(): - pass - - -def _lower__ZL10__hadd2_rn7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hadd2_rn7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hadd2_rn(*a, *b); - return 0; - } - """ - - _ZL10__hadd2_rn7__half2S__1 = declare_device( - "_ZL10__hadd2_rn7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL10__hadd2_rn7__half2S__1_caller(arg_0, arg_1): - return _ZL10__hadd2_rn7__half2S__1(arg_0, arg_1) - - @lower(__hadd2_rn, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hadd2_rn7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hadd2_rn7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL10__hadd2_rn7__half2S__1(shim_stream, shim_obj) - - -def __hsub2_rn(): - pass - - -def _lower__ZL10__hsub2_rn7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hsub2_rn7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hsub2_rn(*a, *b); - return 0; - } - """ - - _ZL10__hsub2_rn7__half2S__1 = declare_device( - "_ZL10__hsub2_rn7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL10__hsub2_rn7__half2S__1_caller(arg_0, arg_1): - return _ZL10__hsub2_rn7__half2S__1(arg_0, arg_1) - - @lower(__hsub2_rn, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hsub2_rn7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hsub2_rn7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL10__hsub2_rn7__half2S__1(shim_stream, shim_obj) - - -def __hmul2_rn(): - pass - - -def _lower__ZL10__hmul2_rn7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hmul2_rn7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hmul2_rn(*a, *b); - return 0; - } - """ - - _ZL10__hmul2_rn7__half2S__1 = declare_device( - "_ZL10__hmul2_rn7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL10__hmul2_rn7__half2S__1_caller(arg_0, arg_1): - return _ZL10__hmul2_rn7__half2S__1(arg_0, arg_1) - - @lower(__hmul2_rn, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hmul2_rn7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hmul2_rn7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL10__hmul2_rn7__half2S__1(shim_stream, shim_obj) - - -def __h2div(): - pass - - -def _lower__ZL7__h2div7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__h2div7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __h2div(*a, *b); - return 0; - } - """ - - _ZL7__h2div7__half2S__1 = declare_device( - "_ZL7__h2div7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__h2div7__half2S__1_caller(arg_0, arg_1): - return _ZL7__h2div7__half2S__1(arg_0, arg_1) - - @lower(__h2div, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__h2div7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__h2div7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL7__h2div7__half2S__1(shim_stream, shim_obj) - - -def __habs2(): - pass - - -def _lower__ZL7__habs27__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__habs27__half2_1(__half2 &retval , __half2* a) { - retval = __habs2(*a); - return 0; - } - """ - - _ZL7__habs27__half2_1 = declare_device( - "_ZL7__habs27__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL7__habs27__half2_1_caller(arg_0): - return _ZL7__habs27__half2_1(arg_0) - - @lower(__habs2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__habs27__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__habs27__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__habs27__half2_1(shim_stream, shim_obj) - - -def __hadd2_sat(): - pass - - -def _lower__ZL11__hadd2_sat7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL11__hadd2_sat7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hadd2_sat(*a, *b); - return 0; - } - """ - - _ZL11__hadd2_sat7__half2S__1 = declare_device( - "_ZL11__hadd2_sat7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL11__hadd2_sat7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hadd2_sat7__half2S__1(arg_0, arg_1) - - @lower(__hadd2_sat, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hadd2_sat7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL11__hadd2_sat7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL11__hadd2_sat7__half2S__1(shim_stream, shim_obj) - - -def __hsub2_sat(): - pass - - -def _lower__ZL11__hsub2_sat7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL11__hsub2_sat7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hsub2_sat(*a, *b); - return 0; - } - """ - - _ZL11__hsub2_sat7__half2S__1 = declare_device( - "_ZL11__hsub2_sat7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL11__hsub2_sat7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hsub2_sat7__half2S__1(arg_0, arg_1) - - @lower(__hsub2_sat, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hsub2_sat7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL11__hsub2_sat7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL11__hsub2_sat7__half2S__1(shim_stream, shim_obj) - - -def __hmul2_sat(): - pass - - -def _lower__ZL11__hmul2_sat7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL11__hmul2_sat7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hmul2_sat(*a, *b); - return 0; - } - """ - - _ZL11__hmul2_sat7__half2S__1 = declare_device( - "_ZL11__hmul2_sat7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL11__hmul2_sat7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hmul2_sat7__half2S__1(arg_0, arg_1) - - @lower(__hmul2_sat, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hmul2_sat7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL11__hmul2_sat7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL11__hmul2_sat7__half2S__1(shim_stream, shim_obj) - - -def __hfma2(): - pass - - -def _lower__ZL7__hfma27__half2S_S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hfma27__half2S_S__1(__half2 &retval , __half2* a, __half2* b, __half2* c) { - retval = __hfma2(*a, *b, *c); - return 0; - } - """ - - _ZL7__hfma27__half2S_S__1 = declare_device( - "_ZL7__hfma27__half2S_S__1", - _type___half2( - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ) - - def _ZL7__hfma27__half2S_S__1_caller(arg_0, arg_1, arg_2): - return _ZL7__hfma27__half2S_S__1(arg_0, arg_1, arg_2) - - @lower(__hfma2, _type___half2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hfma27__half2S_S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hfma27__half2S_S__1_caller, - signature( - _type___half2, - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ptrs, - ) - - -_lower__ZL7__hfma27__half2S_S__1(shim_stream, shim_obj) - - -def __hfma2_sat(): - pass - - -def _lower__ZL11__hfma2_sat7__half2S_S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL11__hfma2_sat7__half2S_S__1(__half2 &retval , __half2* a, __half2* b, __half2* c) { - retval = __hfma2_sat(*a, *b, *c); - return 0; - } - """ - - _ZL11__hfma2_sat7__half2S_S__1 = declare_device( - "_ZL11__hfma2_sat7__half2S_S__1", - _type___half2( - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ) - - def _ZL11__hfma2_sat7__half2S_S__1_caller(arg_0, arg_1, arg_2): - return _ZL11__hfma2_sat7__half2S_S__1(arg_0, arg_1, arg_2) - - @lower(__hfma2_sat, _type___half2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL11__hfma2_sat7__half2S_S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL11__hfma2_sat7__half2S_S__1_caller, - signature( - _type___half2, - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ptrs, - ) - - -_lower__ZL11__hfma2_sat7__half2S_S__1(shim_stream, shim_obj) - - -def __hneg2(): - pass - - -def _lower__ZL7__hneg27__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hneg27__half2_1(__half2 &retval , __half2* a) { - retval = __hneg2(*a); - return 0; - } - """ - - _ZL7__hneg27__half2_1 = declare_device( - "_ZL7__hneg27__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL7__hneg27__half2_1_caller(arg_0): - return _ZL7__hneg27__half2_1(arg_0) - - @lower(__hneg2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hneg27__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hneg27__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__hneg27__half2_1(shim_stream, shim_obj) - - -def __habs(): - pass - - -def _lower__ZL6__habs6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__habs6__half_1(__half &retval , __half* a) { - retval = __habs(*a); - return 0; - } - """ - - _ZL6__habs6__half_1 = declare_device( - "_ZL6__habs6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL6__habs6__half_1_caller(arg_0): - return _ZL6__habs6__half_1(arg_0) - - @lower(__habs, _type___half) - @lower(abs, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__habs6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__habs6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__habs6__half_1(shim_stream, shim_obj) - - -def __hadd(): - pass - - -def _lower__ZL6__hadd6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hadd6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hadd(*a, *b); - return 0; - } - """ - - _ZL6__hadd6__halfS__1 = declare_device( - "_ZL6__hadd6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hadd6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hadd6__halfS__1(arg_0, arg_1) - - @lower(__hadd, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hadd6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hadd6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL6__hadd6__halfS__1(shim_stream, shim_obj) - - -def __hsub(): - pass - - -def _lower__ZL6__hsub6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hsub6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hsub(*a, *b); - return 0; - } - """ - - _ZL6__hsub6__halfS__1 = declare_device( - "_ZL6__hsub6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hsub6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hsub6__halfS__1(arg_0, arg_1) - - @lower(__hsub, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hsub6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hsub6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL6__hsub6__halfS__1(shim_stream, shim_obj) - - -def __hmul(): - pass - - -def _lower__ZL6__hmul6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hmul6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hmul(*a, *b); - return 0; - } - """ - - _ZL6__hmul6__halfS__1 = declare_device( - "_ZL6__hmul6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hmul6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hmul6__halfS__1(arg_0, arg_1) - - @lower(__hmul, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hmul6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hmul6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL6__hmul6__halfS__1(shim_stream, shim_obj) - - -def __hadd_rn(): - pass - - -def _lower__ZL9__hadd_rn6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL9__hadd_rn6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hadd_rn(*a, *b); - return 0; - } - """ - - _ZL9__hadd_rn6__halfS__1 = declare_device( - "_ZL9__hadd_rn6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL9__hadd_rn6__halfS__1_caller(arg_0, arg_1): - return _ZL9__hadd_rn6__halfS__1(arg_0, arg_1) - - @lower(__hadd_rn, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL9__hadd_rn6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL9__hadd_rn6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL9__hadd_rn6__halfS__1(shim_stream, shim_obj) - - -def __hsub_rn(): - pass - - -def _lower__ZL9__hsub_rn6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL9__hsub_rn6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hsub_rn(*a, *b); - return 0; - } - """ - - _ZL9__hsub_rn6__halfS__1 = declare_device( - "_ZL9__hsub_rn6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL9__hsub_rn6__halfS__1_caller(arg_0, arg_1): - return _ZL9__hsub_rn6__halfS__1(arg_0, arg_1) - - @lower(__hsub_rn, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL9__hsub_rn6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL9__hsub_rn6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL9__hsub_rn6__halfS__1(shim_stream, shim_obj) - - -def __hmul_rn(): - pass - - -def _lower__ZL9__hmul_rn6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL9__hmul_rn6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hmul_rn(*a, *b); - return 0; - } - """ - - _ZL9__hmul_rn6__halfS__1 = declare_device( - "_ZL9__hmul_rn6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL9__hmul_rn6__halfS__1_caller(arg_0, arg_1): - return _ZL9__hmul_rn6__halfS__1(arg_0, arg_1) - - @lower(__hmul_rn, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL9__hmul_rn6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL9__hmul_rn6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL9__hmul_rn6__halfS__1(shim_stream, shim_obj) - - -def __hdiv(): - pass - - -def _lower__ZL6__hdiv6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hdiv6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hdiv(*a, *b); - return 0; - } - """ - - _ZL6__hdiv6__halfS__1 = declare_device( - "_ZL6__hdiv6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hdiv6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hdiv6__halfS__1(arg_0, arg_1) - - @lower(__hdiv, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hdiv6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hdiv6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL6__hdiv6__halfS__1(shim_stream, shim_obj) - - -def __hadd_sat(): - pass - - -def _lower__ZL10__hadd_sat6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hadd_sat6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hadd_sat(*a, *b); - return 0; - } - """ - - _ZL10__hadd_sat6__halfS__1 = declare_device( - "_ZL10__hadd_sat6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL10__hadd_sat6__halfS__1_caller(arg_0, arg_1): - return _ZL10__hadd_sat6__halfS__1(arg_0, arg_1) - - @lower(__hadd_sat, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hadd_sat6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hadd_sat6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL10__hadd_sat6__halfS__1(shim_stream, shim_obj) - - -def __hsub_sat(): - pass - - -def _lower__ZL10__hsub_sat6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hsub_sat6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hsub_sat(*a, *b); - return 0; - } - """ - - _ZL10__hsub_sat6__halfS__1 = declare_device( - "_ZL10__hsub_sat6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL10__hsub_sat6__halfS__1_caller(arg_0, arg_1): - return _ZL10__hsub_sat6__halfS__1(arg_0, arg_1) - - @lower(__hsub_sat, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hsub_sat6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hsub_sat6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL10__hsub_sat6__halfS__1(shim_stream, shim_obj) - - -def __hmul_sat(): - pass - - -def _lower__ZL10__hmul_sat6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hmul_sat6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hmul_sat(*a, *b); - return 0; - } - """ - - _ZL10__hmul_sat6__halfS__1 = declare_device( - "_ZL10__hmul_sat6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL10__hmul_sat6__halfS__1_caller(arg_0, arg_1): - return _ZL10__hmul_sat6__halfS__1(arg_0, arg_1) - - @lower(__hmul_sat, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hmul_sat6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hmul_sat6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL10__hmul_sat6__halfS__1(shim_stream, shim_obj) - - -def __hfma(): - pass - - -def _lower__ZL6__hfma6__halfS_S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hfma6__halfS_S__1(__half &retval , __half* a, __half* b, __half* c) { - retval = __hfma(*a, *b, *c); - return 0; - } - """ - - _ZL6__hfma6__halfS_S__1 = declare_device( - "_ZL6__hfma6__halfS_S__1", - _type___half( - CPointer(_type___half), - CPointer(_type___half), - CPointer(_type___half), - ), - ) - - def _ZL6__hfma6__halfS_S__1_caller(arg_0, arg_1, arg_2): - return _ZL6__hfma6__halfS_S__1(arg_0, arg_1, arg_2) - - @lower(__hfma, _type___half, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hfma6__halfS_S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hfma6__halfS_S__1_caller, - signature( - _type___half, - CPointer(_type___half), - CPointer(_type___half), - CPointer(_type___half), - ), - ptrs, - ) - - -_lower__ZL6__hfma6__halfS_S__1(shim_stream, shim_obj) - - -def __hfma_sat(): - pass - - -def _lower__ZL10__hfma_sat6__halfS_S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hfma_sat6__halfS_S__1(__half &retval , __half* a, __half* b, __half* c) { - retval = __hfma_sat(*a, *b, *c); - return 0; - } - """ - - _ZL10__hfma_sat6__halfS_S__1 = declare_device( - "_ZL10__hfma_sat6__halfS_S__1", - _type___half( - CPointer(_type___half), - CPointer(_type___half), - CPointer(_type___half), - ), - ) - - def _ZL10__hfma_sat6__halfS_S__1_caller(arg_0, arg_1, arg_2): - return _ZL10__hfma_sat6__halfS_S__1(arg_0, arg_1, arg_2) - - @lower(__hfma_sat, _type___half, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hfma_sat6__halfS_S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hfma_sat6__halfS_S__1_caller, - signature( - _type___half, - CPointer(_type___half), - CPointer(_type___half), - CPointer(_type___half), - ), - ptrs, - ) - - -_lower__ZL10__hfma_sat6__halfS_S__1(shim_stream, shim_obj) - - -def __hneg(): - pass - - -def _lower__ZL6__hneg6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hneg6__half_1(__half &retval , __half* a) { - retval = __hneg(*a); - return 0; - } - """ - - _ZL6__hneg6__half_1 = declare_device( - "_ZL6__hneg6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL6__hneg6__half_1_caller(arg_0): - return _ZL6__hneg6__half_1(arg_0) - - @lower(__hneg, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hneg6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hneg6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__hneg6__half_1(shim_stream, shim_obj) - - -def __hbeq2(): - pass - - -def _lower__ZL7__hbeq27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hbeq27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbeq2(*a, *b); - return 0; - } - """ - - _ZL7__hbeq27__half2S__1 = declare_device( - "_ZL7__hbeq27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hbeq27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hbeq27__half2S__1(arg_0, arg_1) - - @lower(__hbeq2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hbeq27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hbeq27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__hbeq27__half2S__1(shim_stream, shim_obj) - - -def __hbne2(): - pass - - -def _lower__ZL7__hbne27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hbne27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbne2(*a, *b); - return 0; - } - """ - - _ZL7__hbne27__half2S__1 = declare_device( - "_ZL7__hbne27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hbne27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hbne27__half2S__1(arg_0, arg_1) - - @lower(__hbne2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hbne27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hbne27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__hbne27__half2S__1(shim_stream, shim_obj) - - -def __hble2(): - pass - - -def _lower__ZL7__hble27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hble27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hble2(*a, *b); - return 0; - } - """ - - _ZL7__hble27__half2S__1 = declare_device( - "_ZL7__hble27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hble27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hble27__half2S__1(arg_0, arg_1) - - @lower(__hble2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hble27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hble27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__hble27__half2S__1(shim_stream, shim_obj) - - -def __hbge2(): - pass - - -def _lower__ZL7__hbge27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hbge27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbge2(*a, *b); - return 0; - } - """ - - _ZL7__hbge27__half2S__1 = declare_device( - "_ZL7__hbge27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hbge27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hbge27__half2S__1(arg_0, arg_1) - - @lower(__hbge2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hbge27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hbge27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__hbge27__half2S__1(shim_stream, shim_obj) - - -def __hblt2(): - pass - - -def _lower__ZL7__hblt27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hblt27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hblt2(*a, *b); - return 0; - } - """ - - _ZL7__hblt27__half2S__1 = declare_device( - "_ZL7__hblt27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hblt27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hblt27__half2S__1(arg_0, arg_1) - - @lower(__hblt2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hblt27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hblt27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__hblt27__half2S__1(shim_stream, shim_obj) - - -def __hbgt2(): - pass - - -def _lower__ZL7__hbgt27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7__hbgt27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbgt2(*a, *b); - return 0; - } - """ - - _ZL7__hbgt27__half2S__1 = declare_device( - "_ZL7__hbgt27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL7__hbgt27__half2S__1_caller(arg_0, arg_1): - return _ZL7__hbgt27__half2S__1(arg_0, arg_1) - - @lower(__hbgt2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7__hbgt27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7__hbgt27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7__hbgt27__half2S__1(shim_stream, shim_obj) - - -def __hbequ2(): - pass - - -def _lower__ZL8__hbequ27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hbequ27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbequ2(*a, *b); - return 0; - } - """ - - _ZL8__hbequ27__half2S__1 = declare_device( - "_ZL8__hbequ27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL8__hbequ27__half2S__1_caller(arg_0, arg_1): - return _ZL8__hbequ27__half2S__1(arg_0, arg_1) - - @lower(__hbequ2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hbequ27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hbequ27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL8__hbequ27__half2S__1(shim_stream, shim_obj) - - -def __hbneu2(): - pass - - -def _lower__ZL8__hbneu27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hbneu27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbneu2(*a, *b); - return 0; - } - """ - - _ZL8__hbneu27__half2S__1 = declare_device( - "_ZL8__hbneu27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL8__hbneu27__half2S__1_caller(arg_0, arg_1): - return _ZL8__hbneu27__half2S__1(arg_0, arg_1) - - @lower(__hbneu2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hbneu27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hbneu27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL8__hbneu27__half2S__1(shim_stream, shim_obj) - - -def __hbleu2(): - pass - - -def _lower__ZL8__hbleu27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hbleu27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbleu2(*a, *b); - return 0; - } - """ - - _ZL8__hbleu27__half2S__1 = declare_device( - "_ZL8__hbleu27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL8__hbleu27__half2S__1_caller(arg_0, arg_1): - return _ZL8__hbleu27__half2S__1(arg_0, arg_1) - - @lower(__hbleu2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hbleu27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hbleu27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL8__hbleu27__half2S__1(shim_stream, shim_obj) - - -def __hbgeu2(): - pass - - -def _lower__ZL8__hbgeu27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hbgeu27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbgeu2(*a, *b); - return 0; - } - """ - - _ZL8__hbgeu27__half2S__1 = declare_device( - "_ZL8__hbgeu27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL8__hbgeu27__half2S__1_caller(arg_0, arg_1): - return _ZL8__hbgeu27__half2S__1(arg_0, arg_1) - - @lower(__hbgeu2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hbgeu27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hbgeu27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL8__hbgeu27__half2S__1(shim_stream, shim_obj) - - -def __hbltu2(): - pass - - -def _lower__ZL8__hbltu27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hbltu27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbltu2(*a, *b); - return 0; - } - """ - - _ZL8__hbltu27__half2S__1 = declare_device( - "_ZL8__hbltu27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL8__hbltu27__half2S__1_caller(arg_0, arg_1): - return _ZL8__hbltu27__half2S__1(arg_0, arg_1) - - @lower(__hbltu2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hbltu27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hbltu27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL8__hbltu27__half2S__1(shim_stream, shim_obj) - - -def __hbgtu2(): - pass - - -def _lower__ZL8__hbgtu27__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hbgtu27__half2S__1(bool &retval , __half2* a, __half2* b) { - retval = __hbgtu2(*a, *b); - return 0; - } - """ - - _ZL8__hbgtu27__half2S__1 = declare_device( - "_ZL8__hbgtu27__half2S__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL8__hbgtu27__half2S__1_caller(arg_0, arg_1): - return _ZL8__hbgtu27__half2S__1(arg_0, arg_1) - - @lower(__hbgtu2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hbgtu27__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hbgtu27__half2S__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL8__hbgtu27__half2S__1(shim_stream, shim_obj) - - -def __heq(): - pass - - -def _lower__ZL5__heq6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5__heq6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __heq(*a, *b); - return 0; - } - """ - - _ZL5__heq6__halfS__1 = declare_device( - "_ZL5__heq6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL5__heq6__halfS__1_caller(arg_0, arg_1): - return _ZL5__heq6__halfS__1(arg_0, arg_1) - - @lower(__heq, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__heq6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5__heq6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5__heq6__halfS__1(shim_stream, shim_obj) - - -def __hne(): - pass - - -def _lower__ZL5__hne6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5__hne6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hne(*a, *b); - return 0; - } - """ - - _ZL5__hne6__halfS__1 = declare_device( - "_ZL5__hne6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL5__hne6__halfS__1_caller(arg_0, arg_1): - return _ZL5__hne6__halfS__1(arg_0, arg_1) - - @lower(__hne, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__hne6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5__hne6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5__hne6__halfS__1(shim_stream, shim_obj) - - -def __hle(): - pass - - -def _lower__ZL5__hle6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5__hle6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hle(*a, *b); - return 0; - } - """ - - _ZL5__hle6__halfS__1 = declare_device( - "_ZL5__hle6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL5__hle6__halfS__1_caller(arg_0, arg_1): - return _ZL5__hle6__halfS__1(arg_0, arg_1) - - @lower(__hle, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__hle6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5__hle6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5__hle6__halfS__1(shim_stream, shim_obj) - - -def __hge(): - pass - - -def _lower__ZL5__hge6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5__hge6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hge(*a, *b); - return 0; - } - """ - - _ZL5__hge6__halfS__1 = declare_device( - "_ZL5__hge6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL5__hge6__halfS__1_caller(arg_0, arg_1): - return _ZL5__hge6__halfS__1(arg_0, arg_1) - - @lower(__hge, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__hge6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5__hge6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5__hge6__halfS__1(shim_stream, shim_obj) - - -def __hlt(): - pass - - -def _lower__ZL5__hlt6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5__hlt6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hlt(*a, *b); - return 0; - } - """ - - _ZL5__hlt6__halfS__1 = declare_device( - "_ZL5__hlt6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL5__hlt6__halfS__1_caller(arg_0, arg_1): - return _ZL5__hlt6__halfS__1(arg_0, arg_1) - - @lower(__hlt, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__hlt6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5__hlt6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5__hlt6__halfS__1(shim_stream, shim_obj) - - -def __hgt(): - pass - - -def _lower__ZL5__hgt6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5__hgt6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hgt(*a, *b); - return 0; - } - """ - - _ZL5__hgt6__halfS__1 = declare_device( - "_ZL5__hgt6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL5__hgt6__halfS__1_caller(arg_0, arg_1): - return _ZL5__hgt6__halfS__1(arg_0, arg_1) - - @lower(__hgt, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5__hgt6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5__hgt6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5__hgt6__halfS__1(shim_stream, shim_obj) - - -def __hequ(): - pass - - -def _lower__ZL6__hequ6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hequ6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hequ(*a, *b); - return 0; - } - """ - - _ZL6__hequ6__halfS__1 = declare_device( - "_ZL6__hequ6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hequ6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hequ6__halfS__1(arg_0, arg_1) - - @lower(__hequ, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hequ6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hequ6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__hequ6__halfS__1(shim_stream, shim_obj) - - -def __hneu(): - pass - - -def _lower__ZL6__hneu6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hneu6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hneu(*a, *b); - return 0; - } - """ - - _ZL6__hneu6__halfS__1 = declare_device( - "_ZL6__hneu6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hneu6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hneu6__halfS__1(arg_0, arg_1) - - @lower(__hneu, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hneu6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hneu6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__hneu6__halfS__1(shim_stream, shim_obj) - - -def __hleu(): - pass - - -def _lower__ZL6__hleu6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hleu6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hleu(*a, *b); - return 0; - } - """ - - _ZL6__hleu6__halfS__1 = declare_device( - "_ZL6__hleu6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hleu6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hleu6__halfS__1(arg_0, arg_1) - - @lower(__hleu, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hleu6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hleu6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__hleu6__halfS__1(shim_stream, shim_obj) - - -def __hgeu(): - pass - - -def _lower__ZL6__hgeu6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hgeu6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hgeu(*a, *b); - return 0; - } - """ - - _ZL6__hgeu6__halfS__1 = declare_device( - "_ZL6__hgeu6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hgeu6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hgeu6__halfS__1(arg_0, arg_1) - - @lower(__hgeu, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hgeu6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hgeu6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__hgeu6__halfS__1(shim_stream, shim_obj) - - -def __hltu(): - pass - - -def _lower__ZL6__hltu6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hltu6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hltu(*a, *b); - return 0; - } - """ - - _ZL6__hltu6__halfS__1 = declare_device( - "_ZL6__hltu6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hltu6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hltu6__halfS__1(arg_0, arg_1) - - @lower(__hltu, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hltu6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hltu6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__hltu6__halfS__1(shim_stream, shim_obj) - - -def __hgtu(): - pass - - -def _lower__ZL6__hgtu6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6__hgtu6__halfS__1(bool &retval , __half* a, __half* b) { - retval = __hgtu(*a, *b); - return 0; - } - """ - - _ZL6__hgtu6__halfS__1 = declare_device( - "_ZL6__hgtu6__halfS__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL6__hgtu6__halfS__1_caller(arg_0, arg_1): - return _ZL6__hgtu6__halfS__1(arg_0, arg_1) - - @lower(__hgtu, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6__hgtu6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6__hgtu6__halfS__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6__hgtu6__halfS__1(shim_stream, shim_obj) - - -def __hisnan(): - pass - - -def _lower__ZL8__hisnan6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hisnan6__half_1(bool &retval , __half* a) { - retval = __hisnan(*a); - return 0; - } - """ - - _ZL8__hisnan6__half_1 = declare_device( - "_ZL8__hisnan6__half_1", bool_(CPointer(_type___half)) - ) - - def _ZL8__hisnan6__half_1_caller(arg_0): - return _ZL8__hisnan6__half_1(arg_0) - - @lower(__hisnan, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hisnan6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hisnan6__half_1_caller, - signature(bool_, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL8__hisnan6__half_1(shim_stream, shim_obj) - - -def __hmax_nan(): - pass - - -def _lower__ZL10__hmax_nan6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hmax_nan6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hmax_nan(*a, *b); - return 0; - } - """ - - _ZL10__hmax_nan6__halfS__1 = declare_device( - "_ZL10__hmax_nan6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL10__hmax_nan6__halfS__1_caller(arg_0, arg_1): - return _ZL10__hmax_nan6__halfS__1(arg_0, arg_1) - - @lower(__hmax_nan, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hmax_nan6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hmax_nan6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL10__hmax_nan6__halfS__1(shim_stream, shim_obj) - - -def __hmin_nan(): - pass - - -def _lower__ZL10__hmin_nan6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL10__hmin_nan6__halfS__1(__half &retval , __half* a, __half* b) { - retval = __hmin_nan(*a, *b); - return 0; - } - """ - - _ZL10__hmin_nan6__halfS__1 = declare_device( - "_ZL10__hmin_nan6__halfS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZL10__hmin_nan6__halfS__1_caller(arg_0, arg_1): - return _ZL10__hmin_nan6__halfS__1(arg_0, arg_1) - - @lower(__hmin_nan, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL10__hmin_nan6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL10__hmin_nan6__halfS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZL10__hmin_nan6__halfS__1(shim_stream, shim_obj) - - -def __hfma_relu(): - pass - - -def _lower__ZL11__hfma_relu6__halfS_S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL11__hfma_relu6__halfS_S__1(__half &retval , __half* a, __half* b, __half* c) { - retval = __hfma_relu(*a, *b, *c); - return 0; - } - """ - - _ZL11__hfma_relu6__halfS_S__1 = declare_device( - "_ZL11__hfma_relu6__halfS_S__1", - _type___half( - CPointer(_type___half), - CPointer(_type___half), - CPointer(_type___half), - ), - ) - - def _ZL11__hfma_relu6__halfS_S__1_caller(arg_0, arg_1, arg_2): - return _ZL11__hfma_relu6__halfS_S__1(arg_0, arg_1, arg_2) - - @lower(__hfma_relu, _type___half, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL11__hfma_relu6__halfS_S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL11__hfma_relu6__halfS_S__1_caller, - signature( - _type___half, - CPointer(_type___half), - CPointer(_type___half), - CPointer(_type___half), - ), - ptrs, - ) - - -_lower__ZL11__hfma_relu6__halfS_S__1(shim_stream, shim_obj) - - -def __hmax2_nan(): - pass - - -def _lower__ZL11__hmax2_nan7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL11__hmax2_nan7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hmax2_nan(*a, *b); - return 0; - } - """ - - _ZL11__hmax2_nan7__half2S__1 = declare_device( - "_ZL11__hmax2_nan7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL11__hmax2_nan7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hmax2_nan7__half2S__1(arg_0, arg_1) - - @lower(__hmax2_nan, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hmax2_nan7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL11__hmax2_nan7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL11__hmax2_nan7__half2S__1(shim_stream, shim_obj) - - -def __hmin2_nan(): - pass - - -def _lower__ZL11__hmin2_nan7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL11__hmin2_nan7__half2S__1(__half2 &retval , __half2* a, __half2* b) { - retval = __hmin2_nan(*a, *b); - return 0; - } - """ - - _ZL11__hmin2_nan7__half2S__1 = declare_device( - "_ZL11__hmin2_nan7__half2S__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZL11__hmin2_nan7__half2S__1_caller(arg_0, arg_1): - return _ZL11__hmin2_nan7__half2S__1(arg_0, arg_1) - - @lower(__hmin2_nan, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL11__hmin2_nan7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL11__hmin2_nan7__half2S__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZL11__hmin2_nan7__half2S__1(shim_stream, shim_obj) - - -def __hfma2_relu(): - pass - - -def _lower__ZL12__hfma2_relu7__half2S_S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12__hfma2_relu7__half2S_S__1(__half2 &retval , __half2* a, __half2* b, __half2* c) { - retval = __hfma2_relu(*a, *b, *c); - return 0; - } - """ - - _ZL12__hfma2_relu7__half2S_S__1 = declare_device( - "_ZL12__hfma2_relu7__half2S_S__1", - _type___half2( - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ) - - def _ZL12__hfma2_relu7__half2S_S__1_caller(arg_0, arg_1, arg_2): - return _ZL12__hfma2_relu7__half2S_S__1(arg_0, arg_1, arg_2) - - @lower(__hfma2_relu, _type___half2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key( - "_ZL12__hfma2_relu7__half2S_S__1", shim_raw_str - ) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL12__hfma2_relu7__half2S_S__1_caller, - signature( - _type___half2, - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ptrs, - ) - - -_lower__ZL12__hfma2_relu7__half2S_S__1(shim_stream, shim_obj) - - -def __hcmadd(): - pass - - -def _lower__ZL8__hcmadd7__half2S_S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL8__hcmadd7__half2S_S__1(__half2 &retval , __half2* a, __half2* b, __half2* c) { - retval = __hcmadd(*a, *b, *c); - return 0; - } - """ - - _ZL8__hcmadd7__half2S_S__1 = declare_device( - "_ZL8__hcmadd7__half2S_S__1", - _type___half2( - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ) - - def _ZL8__hcmadd7__half2S_S__1_caller(arg_0, arg_1, arg_2): - return _ZL8__hcmadd7__half2S_S__1(arg_0, arg_1, arg_2) - - @lower(__hcmadd, _type___half2, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL8__hcmadd7__half2S_S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL8__hcmadd7__half2S_S__1_caller, - signature( - _type___half2, - CPointer(_type___half2), - CPointer(_type___half2), - CPointer(_type___half2), - ), - ptrs, - ) - - -_lower__ZL8__hcmadd7__half2S_S__1(shim_stream, shim_obj) - - -def hsqrt(): - pass - - -def _lower__ZL5hsqrt6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5hsqrt6__half_1(__half &retval , __half* a) { - retval = hsqrt(*a); - return 0; - } - """ - - _ZL5hsqrt6__half_1 = declare_device( - "_ZL5hsqrt6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL5hsqrt6__half_1_caller(arg_0): - return _ZL5hsqrt6__half_1(arg_0) - - @lower(hsqrt, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5hsqrt6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5hsqrt6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5hsqrt6__half_1(shim_stream, shim_obj) - - -def hrsqrt(): - pass - - -def _lower__ZL6hrsqrt6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6hrsqrt6__half_1(__half &retval , __half* a) { - retval = hrsqrt(*a); - return 0; - } - """ - - _ZL6hrsqrt6__half_1 = declare_device( - "_ZL6hrsqrt6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL6hrsqrt6__half_1_caller(arg_0): - return _ZL6hrsqrt6__half_1(arg_0) - - @lower(hrsqrt, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6hrsqrt6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6hrsqrt6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6hrsqrt6__half_1(shim_stream, shim_obj) - - -def hrcp(): - pass - - -def _lower__ZL4hrcp6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL4hrcp6__half_1(__half &retval , __half* a) { - retval = hrcp(*a); - return 0; - } - """ - - _ZL4hrcp6__half_1 = declare_device( - "_ZL4hrcp6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL4hrcp6__half_1_caller(arg_0): - return _ZL4hrcp6__half_1(arg_0) - - @lower(hrcp, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL4hrcp6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL4hrcp6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL4hrcp6__half_1(shim_stream, shim_obj) - - -def hlog(): - pass - - -def _lower__ZL4hlog6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL4hlog6__half_1(__half &retval , __half* a) { - retval = hlog(*a); - return 0; - } - """ - - _ZL4hlog6__half_1 = declare_device( - "_ZL4hlog6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL4hlog6__half_1_caller(arg_0): - return _ZL4hlog6__half_1(arg_0) - - @lower(hlog, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL4hlog6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL4hlog6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL4hlog6__half_1(shim_stream, shim_obj) - - -def hlog2(): - pass - - -def _lower__ZL5hlog26__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5hlog26__half_1(__half &retval , __half* a) { - retval = hlog2(*a); - return 0; - } - """ - - _ZL5hlog26__half_1 = declare_device( - "_ZL5hlog26__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL5hlog26__half_1_caller(arg_0): - return _ZL5hlog26__half_1(arg_0) - - @lower(hlog2, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5hlog26__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5hlog26__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5hlog26__half_1(shim_stream, shim_obj) - - -def hlog10(): - pass - - -def _lower__ZL6hlog106__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6hlog106__half_1(__half &retval , __half* a) { - retval = hlog10(*a); - return 0; - } - """ - - _ZL6hlog106__half_1 = declare_device( - "_ZL6hlog106__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL6hlog106__half_1_caller(arg_0): - return _ZL6hlog106__half_1(arg_0) - - @lower(hlog10, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6hlog106__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6hlog106__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6hlog106__half_1(shim_stream, shim_obj) - - -def hexp(): - pass - - -def _lower__ZL4hexp6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL4hexp6__half_1(__half &retval , __half* a) { - retval = hexp(*a); - return 0; - } - """ - - _ZL4hexp6__half_1 = declare_device( - "_ZL4hexp6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL4hexp6__half_1_caller(arg_0): - return _ZL4hexp6__half_1(arg_0) - - @lower(hexp, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL4hexp6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL4hexp6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL4hexp6__half_1(shim_stream, shim_obj) - - -def htanh_approx(): - pass - - -def _lower__ZL12htanh_approx6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL12htanh_approx6__half_1(__half &retval , __half* a) { - retval = htanh_approx(*a); - return 0; - } - """ - - _ZL12htanh_approx6__half_1 = declare_device( - "_ZL12htanh_approx6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL12htanh_approx6__half_1_caller(arg_0): - return _ZL12htanh_approx6__half_1(arg_0) - - @lower(htanh_approx, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL12htanh_approx6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL12htanh_approx6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL12htanh_approx6__half_1(shim_stream, shim_obj) - - -def h2tanh_approx(): - pass - - -def _lower__ZL13h2tanh_approx7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL13h2tanh_approx7__half2_1(__half2 &retval , __half2* a) { - retval = h2tanh_approx(*a); - return 0; - } - """ - - _ZL13h2tanh_approx7__half2_1 = declare_device( - "_ZL13h2tanh_approx7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL13h2tanh_approx7__half2_1_caller(arg_0): - return _ZL13h2tanh_approx7__half2_1(arg_0) - - @lower(h2tanh_approx, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL13h2tanh_approx7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL13h2tanh_approx7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL13h2tanh_approx7__half2_1(shim_stream, shim_obj) - - -def htanh(): - pass - - -def _lower__ZL5htanh6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5htanh6__half_1(__half &retval , __half* a) { - retval = htanh(*a); - return 0; - } - """ - - _ZL5htanh6__half_1 = declare_device( - "_ZL5htanh6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL5htanh6__half_1_caller(arg_0): - return _ZL5htanh6__half_1(arg_0) - - @lower(htanh, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5htanh6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5htanh6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5htanh6__half_1(shim_stream, shim_obj) - - -def h2tanh(): - pass - - -def _lower__ZL6h2tanh7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6h2tanh7__half2_1(__half2 &retval , __half2* a) { - retval = h2tanh(*a); - return 0; - } - """ - - _ZL6h2tanh7__half2_1 = declare_device( - "_ZL6h2tanh7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL6h2tanh7__half2_1_caller(arg_0): - return _ZL6h2tanh7__half2_1(arg_0) - - @lower(h2tanh, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6h2tanh7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6h2tanh7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL6h2tanh7__half2_1(shim_stream, shim_obj) - - -def hexp2(): - pass - - -def _lower__ZL5hexp26__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5hexp26__half_1(__half &retval , __half* a) { - retval = hexp2(*a); - return 0; - } - """ - - _ZL5hexp26__half_1 = declare_device( - "_ZL5hexp26__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL5hexp26__half_1_caller(arg_0): - return _ZL5hexp26__half_1(arg_0) - - @lower(hexp2, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5hexp26__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5hexp26__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL5hexp26__half_1(shim_stream, shim_obj) - - -def hexp10(): - pass - - -def _lower__ZL6hexp106__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6hexp106__half_1(__half &retval , __half* a) { - retval = hexp10(*a); - return 0; - } - """ - - _ZL6hexp106__half_1 = declare_device( - "_ZL6hexp106__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL6hexp106__half_1_caller(arg_0): - return _ZL6hexp106__half_1(arg_0) - - @lower(hexp10, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6hexp106__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6hexp106__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL6hexp106__half_1(shim_stream, shim_obj) - - -def hcos(): - pass - - -def _lower__ZL4hcos6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL4hcos6__half_1(__half &retval , __half* a) { - retval = hcos(*a); - return 0; - } - """ - - _ZL4hcos6__half_1 = declare_device( - "_ZL4hcos6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL4hcos6__half_1_caller(arg_0): - return _ZL4hcos6__half_1(arg_0) - - @lower(hcos, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL4hcos6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL4hcos6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL4hcos6__half_1(shim_stream, shim_obj) - - -def hsin(): - pass - - -def _lower__ZL4hsin6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL4hsin6__half_1(__half &retval , __half* a) { - retval = hsin(*a); - return 0; - } - """ - - _ZL4hsin6__half_1 = declare_device( - "_ZL4hsin6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZL4hsin6__half_1_caller(arg_0): - return _ZL4hsin6__half_1(arg_0) - - @lower(hsin, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL4hsin6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL4hsin6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZL4hsin6__half_1(shim_stream, shim_obj) - - -def h2sqrt(): - pass - - -def _lower__ZL6h2sqrt7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6h2sqrt7__half2_1(__half2 &retval , __half2* a) { - retval = h2sqrt(*a); - return 0; - } - """ - - _ZL6h2sqrt7__half2_1 = declare_device( - "_ZL6h2sqrt7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL6h2sqrt7__half2_1_caller(arg_0): - return _ZL6h2sqrt7__half2_1(arg_0) - - @lower(h2sqrt, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6h2sqrt7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6h2sqrt7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL6h2sqrt7__half2_1(shim_stream, shim_obj) - - -def h2rsqrt(): - pass - - -def _lower__ZL7h2rsqrt7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7h2rsqrt7__half2_1(__half2 &retval , __half2* a) { - retval = h2rsqrt(*a); - return 0; - } - """ - - _ZL7h2rsqrt7__half2_1 = declare_device( - "_ZL7h2rsqrt7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL7h2rsqrt7__half2_1_caller(arg_0): - return _ZL7h2rsqrt7__half2_1(arg_0) - - @lower(h2rsqrt, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7h2rsqrt7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7h2rsqrt7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7h2rsqrt7__half2_1(shim_stream, shim_obj) - - -def h2rcp(): - pass - - -def _lower__ZL5h2rcp7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5h2rcp7__half2_1(__half2 &retval , __half2* a) { - retval = h2rcp(*a); - return 0; - } - """ - - _ZL5h2rcp7__half2_1 = declare_device( - "_ZL5h2rcp7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL5h2rcp7__half2_1_caller(arg_0): - return _ZL5h2rcp7__half2_1(arg_0) - - @lower(h2rcp, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5h2rcp7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5h2rcp7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL5h2rcp7__half2_1(shim_stream, shim_obj) - - -def h2log(): - pass - - -def _lower__ZL5h2log7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5h2log7__half2_1(__half2 &retval , __half2* a) { - retval = h2log(*a); - return 0; - } - """ - - _ZL5h2log7__half2_1 = declare_device( - "_ZL5h2log7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL5h2log7__half2_1_caller(arg_0): - return _ZL5h2log7__half2_1(arg_0) - - @lower(h2log, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5h2log7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5h2log7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL5h2log7__half2_1(shim_stream, shim_obj) - - -def h2log2(): - pass - - -def _lower__ZL6h2log27__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6h2log27__half2_1(__half2 &retval , __half2* a) { - retval = h2log2(*a); - return 0; - } - """ - - _ZL6h2log27__half2_1 = declare_device( - "_ZL6h2log27__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL6h2log27__half2_1_caller(arg_0): - return _ZL6h2log27__half2_1(arg_0) - - @lower(h2log2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6h2log27__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6h2log27__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL6h2log27__half2_1(shim_stream, shim_obj) - - -def h2log10(): - pass - - -def _lower__ZL7h2log107__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7h2log107__half2_1(__half2 &retval , __half2* a) { - retval = h2log10(*a); - return 0; - } - """ - - _ZL7h2log107__half2_1 = declare_device( - "_ZL7h2log107__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL7h2log107__half2_1_caller(arg_0): - return _ZL7h2log107__half2_1(arg_0) - - @lower(h2log10, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7h2log107__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7h2log107__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7h2log107__half2_1(shim_stream, shim_obj) - - -def h2exp(): - pass - - -def _lower__ZL5h2exp7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5h2exp7__half2_1(__half2 &retval , __half2* a) { - retval = h2exp(*a); - return 0; - } - """ - - _ZL5h2exp7__half2_1 = declare_device( - "_ZL5h2exp7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL5h2exp7__half2_1_caller(arg_0): - return _ZL5h2exp7__half2_1(arg_0) - - @lower(h2exp, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5h2exp7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5h2exp7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL5h2exp7__half2_1(shim_stream, shim_obj) - - -def h2exp2(): - pass - - -def _lower__ZL6h2exp27__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL6h2exp27__half2_1(__half2 &retval , __half2* a) { - retval = h2exp2(*a); - return 0; - } - """ - - _ZL6h2exp27__half2_1 = declare_device( - "_ZL6h2exp27__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL6h2exp27__half2_1_caller(arg_0): - return _ZL6h2exp27__half2_1(arg_0) - - @lower(h2exp2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL6h2exp27__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL6h2exp27__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL6h2exp27__half2_1(shim_stream, shim_obj) - - -def h2exp10(): - pass - - -def _lower__ZL7h2exp107__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL7h2exp107__half2_1(__half2 &retval , __half2* a) { - retval = h2exp10(*a); - return 0; - } - """ - - _ZL7h2exp107__half2_1 = declare_device( - "_ZL7h2exp107__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL7h2exp107__half2_1_caller(arg_0): - return _ZL7h2exp107__half2_1(arg_0) - - @lower(h2exp10, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL7h2exp107__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL7h2exp107__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL7h2exp107__half2_1(shim_stream, shim_obj) - - -def h2cos(): - pass - - -def _lower__ZL5h2cos7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5h2cos7__half2_1(__half2 &retval , __half2* a) { - retval = h2cos(*a); - return 0; - } - """ - - _ZL5h2cos7__half2_1 = declare_device( - "_ZL5h2cos7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL5h2cos7__half2_1_caller(arg_0): - return _ZL5h2cos7__half2_1(arg_0) - - @lower(h2cos, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5h2cos7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5h2cos7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL5h2cos7__half2_1(shim_stream, shim_obj) - - -def h2sin(): - pass - - -def _lower__ZL5h2sin7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL5h2sin7__half2_1(__half2 &retval , __half2* a) { - retval = h2sin(*a); - return 0; - } - """ - - _ZL5h2sin7__half2_1 = declare_device( - "_ZL5h2sin7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZL5h2sin7__half2_1_caller(arg_0): - return _ZL5h2sin7__half2_1(arg_0) - - @lower(h2sin, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL5h2sin7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL5h2sin7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZL5h2sin7__half2_1(shim_stream, shim_obj) - - -def atomicAdd(): - pass - - -def _lower__ZL9atomicAddP7__half2S__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL9atomicAddP7__half2S__1(__half2 &retval , __half2 ** address, __half2* val) { - retval = atomicAdd(*address, *val); - return 0; - } - """ - - _ZL9atomicAddP7__half2S__1 = declare_device( - "_ZL9atomicAddP7__half2S__1", - _type___half2( - CPointer(CPointer(_type___half2)), CPointer(_type___half2) - ), - ) - - def _ZL9atomicAddP7__half2S__1_caller(arg_0, arg_1): - return _ZL9atomicAddP7__half2S__1(arg_0, arg_1) - - @lower(atomicAdd, CPointer(_type___half2), _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL9atomicAddP7__half2S__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL9atomicAddP7__half2S__1_caller, - signature( - _type___half2, - CPointer(CPointer(_type___half2)), - CPointer(_type___half2), - ), - ptrs, - ) - - -_lower__ZL9atomicAddP7__half2S__1(shim_stream, shim_obj) - - -def _lower__ZL9atomicAddP6__halfS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZL9atomicAddP6__halfS__1(__half &retval , __half ** address, __half* val) { - retval = atomicAdd(*address, *val); - return 0; - } - """ - - _ZL9atomicAddP6__halfS__1 = declare_device( - "_ZL9atomicAddP6__halfS__1", - _type___half(CPointer(CPointer(_type___half)), CPointer(_type___half)), - ) - - def _ZL9atomicAddP6__halfS__1_caller(arg_0, arg_1): - return _ZL9atomicAddP6__halfS__1(arg_0, arg_1) - - @lower(atomicAdd, CPointer(_type___half), _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZL9atomicAddP6__halfS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZL9atomicAddP6__halfS__1_caller, - signature( - _type___half, - CPointer(CPointer(_type___half)), - CPointer(_type___half), - ), - ptrs, - ) - - -_lower__ZL9atomicAddP6__halfS__1(shim_stream, shim_obj) - - -def _lower__ZplRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZplRK6__halfS1__1(__half &retval , __half* lh, __half* rh) { - retval = operator+(*lh, *rh); - return 0; - } - """ - - _ZplRK6__halfS1__1 = declare_device( - "_ZplRK6__halfS1__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZplRK6__halfS1__1_caller(arg_0, arg_1): - return _ZplRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.add, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZplRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZplRK6__halfS1__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZplRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZmiRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmiRK6__halfS1__1(__half &retval , __half* lh, __half* rh) { - retval = operator-(*lh, *rh); - return 0; - } - """ - - _ZmiRK6__halfS1__1 = declare_device( - "_ZmiRK6__halfS1__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZmiRK6__halfS1__1_caller(arg_0, arg_1): - return _ZmiRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.sub, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmiRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmiRK6__halfS1__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZmiRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZmlRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmlRK6__halfS1__1(__half &retval , __half* lh, __half* rh) { - retval = operator*(*lh, *rh); - return 0; - } - """ - - _ZmlRK6__halfS1__1 = declare_device( - "_ZmlRK6__halfS1__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZmlRK6__halfS1__1_caller(arg_0, arg_1): - return _ZmlRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.mul, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmlRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmlRK6__halfS1__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZmlRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZdvRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZdvRK6__halfS1__1(__half &retval , __half* lh, __half* rh) { - retval = operator/(*lh, *rh); - return 0; - } - """ - - _ZdvRK6__halfS1__1 = declare_device( - "_ZdvRK6__halfS1__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZdvRK6__halfS1__1_caller(arg_0, arg_1): - return _ZdvRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.truediv, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZdvRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZdvRK6__halfS1__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZdvRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZpLR6__halfRKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZpLR6__halfRKS__1(__half &retval , __half* lh, __half* rh) { - retval = operator+=(*lh, *rh); - return 0; - } - """ - - _ZpLR6__halfRKS__1 = declare_device( - "_ZpLR6__halfRKS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZpLR6__halfRKS__1_caller(arg_0, arg_1): - return _ZpLR6__halfRKS__1(arg_0, arg_1) - - @lower(operator.iadd, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZpLR6__halfRKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZpLR6__halfRKS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZpLR6__halfRKS__1(shim_stream, shim_obj) - - -def _lower__ZmIR6__halfRKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmIR6__halfRKS__1(__half &retval , __half* lh, __half* rh) { - retval = operator-=(*lh, *rh); - return 0; - } - """ - - _ZmIR6__halfRKS__1 = declare_device( - "_ZmIR6__halfRKS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZmIR6__halfRKS__1_caller(arg_0, arg_1): - return _ZmIR6__halfRKS__1(arg_0, arg_1) - - @lower(operator.isub, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmIR6__halfRKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmIR6__halfRKS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZmIR6__halfRKS__1(shim_stream, shim_obj) - - -def _lower__ZmLR6__halfRKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmLR6__halfRKS__1(__half &retval , __half* lh, __half* rh) { - retval = operator*=(*lh, *rh); - return 0; - } - """ - - _ZmLR6__halfRKS__1 = declare_device( - "_ZmLR6__halfRKS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZmLR6__halfRKS__1_caller(arg_0, arg_1): - return _ZmLR6__halfRKS__1(arg_0, arg_1) - - @lower(operator.imul, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmLR6__halfRKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmLR6__halfRKS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZmLR6__halfRKS__1(shim_stream, shim_obj) - - -def _lower__ZdVR6__halfRKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZdVR6__halfRKS__1(__half &retval , __half* lh, __half* rh) { - retval = operator/=(*lh, *rh); - return 0; - } - """ - - _ZdVR6__halfRKS__1 = declare_device( - "_ZdVR6__halfRKS__1", - _type___half(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZdVR6__halfRKS__1_caller(arg_0, arg_1): - return _ZdVR6__halfRKS__1(arg_0, arg_1) - - @lower(operator.itruediv, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZdVR6__halfRKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZdVR6__halfRKS__1_caller, - signature( - _type___half, CPointer(_type___half), CPointer(_type___half) - ), - ptrs, - ) - - -_lower__ZdVR6__halfRKS__1(shim_stream, shim_obj) - - -def _lower__ZpsRK6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZpsRK6__half_1(__half &retval , __half* h) { - retval = operator+(*h); - return 0; - } - """ - - _ZpsRK6__half_1 = declare_device( - "_ZpsRK6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZpsRK6__half_1_caller(arg_0): - return _ZpsRK6__half_1(arg_0) - - @lower(operator.pos, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZpsRK6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZpsRK6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZpsRK6__half_1(shim_stream, shim_obj) - - -def _lower__ZngRK6__half_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZngRK6__half_1(__half &retval , __half* h) { - retval = operator-(*h); - return 0; - } - """ - - _ZngRK6__half_1 = declare_device( - "_ZngRK6__half_1", _type___half(CPointer(_type___half)) - ) - - def _ZngRK6__half_1_caller(arg_0): - return _ZngRK6__half_1(arg_0) - - @lower(operator.neg, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZngRK6__half_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZngRK6__half_1_caller, - signature(_type___half, CPointer(_type___half)), - ptrs, - ) - - -_lower__ZngRK6__half_1(shim_stream, shim_obj) - - -def _lower__ZeqRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZeqRK6__halfS1__1(bool &retval , __half* lh, __half* rh) { - retval = operator==(*lh, *rh); - return 0; - } - """ - - _ZeqRK6__halfS1__1 = declare_device( - "_ZeqRK6__halfS1__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZeqRK6__halfS1__1_caller(arg_0, arg_1): - return _ZeqRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.eq, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZeqRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZeqRK6__halfS1__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZeqRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZneRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZneRK6__halfS1__1(bool &retval , __half* lh, __half* rh) { - retval = operator!=(*lh, *rh); - return 0; - } - """ - - _ZneRK6__halfS1__1 = declare_device( - "_ZneRK6__halfS1__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZneRK6__halfS1__1_caller(arg_0, arg_1): - return _ZneRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.ne, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZneRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZneRK6__halfS1__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZneRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZgtRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZgtRK6__halfS1__1(bool &retval , __half* lh, __half* rh) { - retval = operator>(*lh, *rh); - return 0; - } - """ - - _ZgtRK6__halfS1__1 = declare_device( - "_ZgtRK6__halfS1__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZgtRK6__halfS1__1_caller(arg_0, arg_1): - return _ZgtRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.gt, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZgtRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZgtRK6__halfS1__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZgtRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZltRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZltRK6__halfS1__1(bool &retval , __half* lh, __half* rh) { - retval = operator<(*lh, *rh); - return 0; - } - """ - - _ZltRK6__halfS1__1 = declare_device( - "_ZltRK6__halfS1__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZltRK6__halfS1__1_caller(arg_0, arg_1): - return _ZltRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.lt, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZltRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZltRK6__halfS1__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZltRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZgeRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZgeRK6__halfS1__1(bool &retval , __half* lh, __half* rh) { - retval = operator>=(*lh, *rh); - return 0; - } - """ - - _ZgeRK6__halfS1__1 = declare_device( - "_ZgeRK6__halfS1__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZgeRK6__halfS1__1_caller(arg_0, arg_1): - return _ZgeRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.ge, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZgeRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZgeRK6__halfS1__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZgeRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZleRK6__halfS1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZleRK6__halfS1__1(bool &retval , __half* lh, __half* rh) { - retval = operator<=(*lh, *rh); - return 0; - } - """ - - _ZleRK6__halfS1__1 = declare_device( - "_ZleRK6__halfS1__1", - bool_(CPointer(_type___half), CPointer(_type___half)), - ) - - def _ZleRK6__halfS1__1_caller(arg_0, arg_1): - return _ZleRK6__halfS1__1(arg_0, arg_1) - - @lower(operator.le, _type___half, _type___half) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZleRK6__halfS1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZleRK6__halfS1__1_caller, - signature(bool_, CPointer(_type___half), CPointer(_type___half)), - ptrs, - ) - - -_lower__ZleRK6__halfS1__1(shim_stream, shim_obj) - - -def _lower__ZplRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZplRK7__half2S1__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator+(*lh, *rh); - return 0; - } - """ - - _ZplRK7__half2S1__1 = declare_device( - "_ZplRK7__half2S1__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZplRK7__half2S1__1_caller(arg_0, arg_1): - return _ZplRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.add, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZplRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZplRK7__half2S1__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZplRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZmiRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmiRK7__half2S1__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator-(*lh, *rh); - return 0; - } - """ - - _ZmiRK7__half2S1__1 = declare_device( - "_ZmiRK7__half2S1__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZmiRK7__half2S1__1_caller(arg_0, arg_1): - return _ZmiRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.sub, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmiRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmiRK7__half2S1__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZmiRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZmlRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmlRK7__half2S1__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator*(*lh, *rh); - return 0; - } - """ - - _ZmlRK7__half2S1__1 = declare_device( - "_ZmlRK7__half2S1__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZmlRK7__half2S1__1_caller(arg_0, arg_1): - return _ZmlRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.mul, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmlRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmlRK7__half2S1__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZmlRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZdvRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZdvRK7__half2S1__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator/(*lh, *rh); - return 0; - } - """ - - _ZdvRK7__half2S1__1 = declare_device( - "_ZdvRK7__half2S1__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZdvRK7__half2S1__1_caller(arg_0, arg_1): - return _ZdvRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.truediv, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZdvRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZdvRK7__half2S1__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZdvRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZpLR7__half2RKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZpLR7__half2RKS__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator+=(*lh, *rh); - return 0; - } - """ - - _ZpLR7__half2RKS__1 = declare_device( - "_ZpLR7__half2RKS__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZpLR7__half2RKS__1_caller(arg_0, arg_1): - return _ZpLR7__half2RKS__1(arg_0, arg_1) - - @lower(operator.iadd, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZpLR7__half2RKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZpLR7__half2RKS__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZpLR7__half2RKS__1(shim_stream, shim_obj) - - -def _lower__ZmIR7__half2RKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmIR7__half2RKS__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator-=(*lh, *rh); - return 0; - } - """ - - _ZmIR7__half2RKS__1 = declare_device( - "_ZmIR7__half2RKS__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZmIR7__half2RKS__1_caller(arg_0, arg_1): - return _ZmIR7__half2RKS__1(arg_0, arg_1) - - @lower(operator.isub, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmIR7__half2RKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmIR7__half2RKS__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZmIR7__half2RKS__1(shim_stream, shim_obj) - - -def _lower__ZmLR7__half2RKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZmLR7__half2RKS__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator*=(*lh, *rh); - return 0; - } - """ - - _ZmLR7__half2RKS__1 = declare_device( - "_ZmLR7__half2RKS__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZmLR7__half2RKS__1_caller(arg_0, arg_1): - return _ZmLR7__half2RKS__1(arg_0, arg_1) - - @lower(operator.imul, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZmLR7__half2RKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZmLR7__half2RKS__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZmLR7__half2RKS__1(shim_stream, shim_obj) - - -def _lower__ZdVR7__half2RKS__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZdVR7__half2RKS__1(__half2 &retval , __half2* lh, __half2* rh) { - retval = operator/=(*lh, *rh); - return 0; - } - """ - - _ZdVR7__half2RKS__1 = declare_device( - "_ZdVR7__half2RKS__1", - _type___half2(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZdVR7__half2RKS__1_caller(arg_0, arg_1): - return _ZdVR7__half2RKS__1(arg_0, arg_1) - - @lower(operator.itruediv, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZdVR7__half2RKS__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZdVR7__half2RKS__1_caller, - signature( - _type___half2, CPointer(_type___half2), CPointer(_type___half2) - ), - ptrs, - ) - - -_lower__ZdVR7__half2RKS__1(shim_stream, shim_obj) - - -def _lower__ZpsRK7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZpsRK7__half2_1(__half2 &retval , __half2* h) { - retval = operator+(*h); - return 0; - } - """ - - _ZpsRK7__half2_1 = declare_device( - "_ZpsRK7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZpsRK7__half2_1_caller(arg_0): - return _ZpsRK7__half2_1(arg_0) - - @lower(operator.pos, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZpsRK7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZpsRK7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZpsRK7__half2_1(shim_stream, shim_obj) - - -def _lower__ZngRK7__half2_1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZngRK7__half2_1(__half2 &retval , __half2* h) { - retval = operator-(*h); - return 0; - } - """ - - _ZngRK7__half2_1 = declare_device( - "_ZngRK7__half2_1", _type___half2(CPointer(_type___half2)) - ) - - def _ZngRK7__half2_1_caller(arg_0): - return _ZngRK7__half2_1(arg_0) - - @lower(operator.neg, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZngRK7__half2_1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZngRK7__half2_1_caller, - signature(_type___half2, CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZngRK7__half2_1(shim_stream, shim_obj) - - -def _lower__ZeqRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZeqRK7__half2S1__1(bool &retval , __half2* lh, __half2* rh) { - retval = operator==(*lh, *rh); - return 0; - } - """ - - _ZeqRK7__half2S1__1 = declare_device( - "_ZeqRK7__half2S1__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZeqRK7__half2S1__1_caller(arg_0, arg_1): - return _ZeqRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.eq, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZeqRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZeqRK7__half2S1__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZeqRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZneRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZneRK7__half2S1__1(bool &retval , __half2* lh, __half2* rh) { - retval = operator!=(*lh, *rh); - return 0; - } - """ - - _ZneRK7__half2S1__1 = declare_device( - "_ZneRK7__half2S1__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZneRK7__half2S1__1_caller(arg_0, arg_1): - return _ZneRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.ne, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZneRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZneRK7__half2S1__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZneRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZgtRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZgtRK7__half2S1__1(bool &retval , __half2* lh, __half2* rh) { - retval = operator>(*lh, *rh); - return 0; - } - """ - - _ZgtRK7__half2S1__1 = declare_device( - "_ZgtRK7__half2S1__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZgtRK7__half2S1__1_caller(arg_0, arg_1): - return _ZgtRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.gt, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZgtRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZgtRK7__half2S1__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZgtRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZltRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZltRK7__half2S1__1(bool &retval , __half2* lh, __half2* rh) { - retval = operator<(*lh, *rh); - return 0; - } - """ - - _ZltRK7__half2S1__1 = declare_device( - "_ZltRK7__half2S1__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZltRK7__half2S1__1_caller(arg_0, arg_1): - return _ZltRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.lt, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZltRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZltRK7__half2S1__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZltRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZgeRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZgeRK7__half2S1__1(bool &retval , __half2* lh, __half2* rh) { - retval = operator>=(*lh, *rh); - return 0; - } - """ - - _ZgeRK7__half2S1__1 = declare_device( - "_ZgeRK7__half2S1__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZgeRK7__half2S1__1_caller(arg_0, arg_1): - return _ZgeRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.ge, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZgeRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZgeRK7__half2S1__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZgeRK7__half2S1__1(shim_stream, shim_obj) - - -def _lower__ZleRK7__half2S1__1(shim_stream, shim_obj): - shim_raw_str = """ - extern "C" __device__ int - _ZleRK7__half2S1__1(bool &retval , __half2* lh, __half2* rh) { - retval = operator<=(*lh, *rh); - return 0; - } - """ - - _ZleRK7__half2S1__1 = declare_device( - "_ZleRK7__half2S1__1", - bool_(CPointer(_type___half2), CPointer(_type___half2)), - ) - - def _ZleRK7__half2S1__1_caller(arg_0, arg_1): - return _ZleRK7__half2S1__1(arg_0, arg_1) - - @lower(operator.le, _type___half2, _type___half2) - def impl(context, builder, sig, args): - context.active_code_library.add_linking_file(shim_obj) - shim_stream.write_with_key("_ZleRK7__half2S1__1", shim_raw_str) - ptrs = [builder.alloca(context.get_value_type(arg)) for arg in sig.args] - for ptr, ty, arg in zip(ptrs, sig.args, args): - builder.store(arg, ptr, align=getattr(ty, "alignof_", None)) - - return context.compile_internal( - builder, - _ZleRK7__half2S1__1_caller, - signature(bool_, CPointer(_type___half2), CPointer(_type___half2)), - ptrs, - ) - - -_lower__ZleRK7__half2S1__1(shim_stream, shim_obj) - - -@register -class _typing___double2half(ConcreteTemplate): - key = globals()["__double2half"] - cases = [signature(_type___half, float64)] - - -register_global(__double2half, types.Function(_typing___double2half)) - - -@register -class _typing___float2half(ConcreteTemplate): - key = globals()["__float2half"] - cases = [signature(_type___half, float32)] - - -register_global(__float2half, types.Function(_typing___float2half)) - - -@register -class _typing___float2half_rn(ConcreteTemplate): - key = globals()["__float2half_rn"] - cases = [signature(_type___half, float32)] - - -register_global(__float2half_rn, types.Function(_typing___float2half_rn)) - - -@register -class _typing___float2half_rz(ConcreteTemplate): - key = globals()["__float2half_rz"] - cases = [signature(_type___half, float32)] - - -register_global(__float2half_rz, types.Function(_typing___float2half_rz)) - - -@register -class _typing___float2half_rd(ConcreteTemplate): - key = globals()["__float2half_rd"] - cases = [signature(_type___half, float32)] - - -register_global(__float2half_rd, types.Function(_typing___float2half_rd)) - - -@register -class _typing___float2half_ru(ConcreteTemplate): - key = globals()["__float2half_ru"] - cases = [signature(_type___half, float32)] - - -register_global(__float2half_ru, types.Function(_typing___float2half_ru)) - - -@register -class _typing___half2float(ConcreteTemplate): - key = globals()["__half2float"] - cases = [signature(float32, _type___half)] - - -register_global(__half2float, types.Function(_typing___half2float)) - - -@register -class _typing___float2half2_rn(ConcreteTemplate): - key = globals()["__float2half2_rn"] - cases = [signature(_type___half2, float32)] - - -register_global(__float2half2_rn, types.Function(_typing___float2half2_rn)) - - -@register -class _typing___floats2half2_rn(ConcreteTemplate): - key = globals()["__floats2half2_rn"] - cases = [signature(_type___half2, float32, float32)] - - -register_global(__floats2half2_rn, types.Function(_typing___floats2half2_rn)) - - -@register -class _typing___low2float(ConcreteTemplate): - key = globals()["__low2float"] - cases = [signature(float32, _type___half2)] - - -register_global(__low2float, types.Function(_typing___low2float)) - - -@register -class _typing___high2float(ConcreteTemplate): - key = globals()["__high2float"] - cases = [signature(float32, _type___half2)] - - -register_global(__high2float, types.Function(_typing___high2float)) - - -@register -class _typing___half2char_rz(ConcreteTemplate): - key = globals()["__half2char_rz"] - cases = [signature(int8, _type___half)] - - -register_global(__half2char_rz, types.Function(_typing___half2char_rz)) - - -@register -class _typing___half2uchar_rz(ConcreteTemplate): - key = globals()["__half2uchar_rz"] - cases = [signature(uint8, _type___half)] - - -register_global(__half2uchar_rz, types.Function(_typing___half2uchar_rz)) - - -@register -class _typing___half2short_rz(ConcreteTemplate): - key = globals()["__half2short_rz"] - cases = [signature(int16, _type___half)] - - -register_global(__half2short_rz, types.Function(_typing___half2short_rz)) - - -@register -class _typing___half2ushort_rz(ConcreteTemplate): - key = globals()["__half2ushort_rz"] - cases = [signature(uint16, _type___half)] - - -register_global(__half2ushort_rz, types.Function(_typing___half2ushort_rz)) - - -@register -class _typing___half2int_rz(ConcreteTemplate): - key = globals()["__half2int_rz"] - cases = [signature(int32, _type___half)] - - -register_global(__half2int_rz, types.Function(_typing___half2int_rz)) - - -@register -class _typing___half2uint_rz(ConcreteTemplate): - key = globals()["__half2uint_rz"] - cases = [signature(uint32, _type___half)] - - -register_global(__half2uint_rz, types.Function(_typing___half2uint_rz)) - - -@register -class _typing___half2ll_rz(ConcreteTemplate): - key = globals()["__half2ll_rz"] - cases = [signature(int64, _type___half)] - - -register_global(__half2ll_rz, types.Function(_typing___half2ll_rz)) - - -@register -class _typing___half2ull_rz(ConcreteTemplate): - key = globals()["__half2ull_rz"] - cases = [signature(uint64, _type___half)] - - -register_global(__half2ull_rz, types.Function(_typing___half2ull_rz)) - - -@register -class _typing_make_half2(ConcreteTemplate): - key = globals()["make_half2"] - cases = [signature(_type___half2, _type___half, _type___half)] - - -register_global(make_half2, types.Function(_typing_make_half2)) - - -@register -class _typing___float22half2_rn(ConcreteTemplate): - key = globals()["__float22half2_rn"] - cases = [signature(_type___half2, float32x2)] - - -register_global(__float22half2_rn, types.Function(_typing___float22half2_rn)) - - -@register -class _typing___half22float2(ConcreteTemplate): - key = globals()["__half22float2"] - cases = [signature(float32x2, _type___half2)] - - -register_global(__half22float2, types.Function(_typing___half22float2)) - - -@register -class _typing___half2int_rn(ConcreteTemplate): - key = globals()["__half2int_rn"] - cases = [signature(int32, _type___half)] - - -register_global(__half2int_rn, types.Function(_typing___half2int_rn)) - - -@register -class _typing___half2int_rd(ConcreteTemplate): - key = globals()["__half2int_rd"] - cases = [signature(int32, _type___half)] - - -register_global(__half2int_rd, types.Function(_typing___half2int_rd)) - - -@register -class _typing___half2int_ru(ConcreteTemplate): - key = globals()["__half2int_ru"] - cases = [signature(int32, _type___half)] - - -register_global(__half2int_ru, types.Function(_typing___half2int_ru)) - - -@register -class _typing___int2half_rn(ConcreteTemplate): - key = globals()["__int2half_rn"] - cases = [signature(_type___half, int32)] - - -register_global(__int2half_rn, types.Function(_typing___int2half_rn)) - - -@register -class _typing___int2half_rz(ConcreteTemplate): - key = globals()["__int2half_rz"] - cases = [signature(_type___half, int32)] - - -register_global(__int2half_rz, types.Function(_typing___int2half_rz)) - - -@register -class _typing___int2half_rd(ConcreteTemplate): - key = globals()["__int2half_rd"] - cases = [signature(_type___half, int32)] - - -register_global(__int2half_rd, types.Function(_typing___int2half_rd)) - - -@register -class _typing___int2half_ru(ConcreteTemplate): - key = globals()["__int2half_ru"] - cases = [signature(_type___half, int32)] - - -register_global(__int2half_ru, types.Function(_typing___int2half_ru)) - - -@register -class _typing___half2short_rn(ConcreteTemplate): - key = globals()["__half2short_rn"] - cases = [signature(int16, _type___half)] - - -register_global(__half2short_rn, types.Function(_typing___half2short_rn)) - - -@register -class _typing___half2short_rd(ConcreteTemplate): - key = globals()["__half2short_rd"] - cases = [signature(int16, _type___half)] - - -register_global(__half2short_rd, types.Function(_typing___half2short_rd)) - - -@register -class _typing___half2short_ru(ConcreteTemplate): - key = globals()["__half2short_ru"] - cases = [signature(int16, _type___half)] - - -register_global(__half2short_ru, types.Function(_typing___half2short_ru)) - - -@register -class _typing___short2half_rn(ConcreteTemplate): - key = globals()["__short2half_rn"] - cases = [signature(_type___half, int16)] - - -register_global(__short2half_rn, types.Function(_typing___short2half_rn)) - - -@register -class _typing___short2half_rz(ConcreteTemplate): - key = globals()["__short2half_rz"] - cases = [signature(_type___half, int16)] - - -register_global(__short2half_rz, types.Function(_typing___short2half_rz)) - - -@register -class _typing___short2half_rd(ConcreteTemplate): - key = globals()["__short2half_rd"] - cases = [signature(_type___half, int16)] - - -register_global(__short2half_rd, types.Function(_typing___short2half_rd)) - - -@register -class _typing___short2half_ru(ConcreteTemplate): - key = globals()["__short2half_ru"] - cases = [signature(_type___half, int16)] - - -register_global(__short2half_ru, types.Function(_typing___short2half_ru)) - - -@register -class _typing___half2uint_rn(ConcreteTemplate): - key = globals()["__half2uint_rn"] - cases = [signature(uint32, _type___half)] - - -register_global(__half2uint_rn, types.Function(_typing___half2uint_rn)) - - -@register -class _typing___half2uint_rd(ConcreteTemplate): - key = globals()["__half2uint_rd"] - cases = [signature(uint32, _type___half)] - - -register_global(__half2uint_rd, types.Function(_typing___half2uint_rd)) - - -@register -class _typing___half2uint_ru(ConcreteTemplate): - key = globals()["__half2uint_ru"] - cases = [signature(uint32, _type___half)] - - -register_global(__half2uint_ru, types.Function(_typing___half2uint_ru)) - - -@register -class _typing___uint2half_rn(ConcreteTemplate): - key = globals()["__uint2half_rn"] - cases = [signature(_type___half, uint32)] - - -register_global(__uint2half_rn, types.Function(_typing___uint2half_rn)) - - -@register -class _typing___uint2half_rz(ConcreteTemplate): - key = globals()["__uint2half_rz"] - cases = [signature(_type___half, uint32)] - - -register_global(__uint2half_rz, types.Function(_typing___uint2half_rz)) - - -@register -class _typing___uint2half_rd(ConcreteTemplate): - key = globals()["__uint2half_rd"] - cases = [signature(_type___half, uint32)] - - -register_global(__uint2half_rd, types.Function(_typing___uint2half_rd)) - - -@register -class _typing___uint2half_ru(ConcreteTemplate): - key = globals()["__uint2half_ru"] - cases = [signature(_type___half, uint32)] - - -register_global(__uint2half_ru, types.Function(_typing___uint2half_ru)) - - -@register -class _typing___half2ushort_rn(ConcreteTemplate): - key = globals()["__half2ushort_rn"] - cases = [signature(uint16, _type___half)] - - -register_global(__half2ushort_rn, types.Function(_typing___half2ushort_rn)) - - -@register -class _typing___half2ushort_rd(ConcreteTemplate): - key = globals()["__half2ushort_rd"] - cases = [signature(uint16, _type___half)] - - -register_global(__half2ushort_rd, types.Function(_typing___half2ushort_rd)) - - -@register -class _typing___half2ushort_ru(ConcreteTemplate): - key = globals()["__half2ushort_ru"] - cases = [signature(uint16, _type___half)] - - -register_global(__half2ushort_ru, types.Function(_typing___half2ushort_ru)) - - -@register -class _typing___ushort2half_rn(ConcreteTemplate): - key = globals()["__ushort2half_rn"] - cases = [signature(_type___half, uint16)] - - -register_global(__ushort2half_rn, types.Function(_typing___ushort2half_rn)) - - -@register -class _typing___ushort2half_rz(ConcreteTemplate): - key = globals()["__ushort2half_rz"] - cases = [signature(_type___half, uint16)] - - -register_global(__ushort2half_rz, types.Function(_typing___ushort2half_rz)) - - -@register -class _typing___ushort2half_rd(ConcreteTemplate): - key = globals()["__ushort2half_rd"] - cases = [signature(_type___half, uint16)] - - -register_global(__ushort2half_rd, types.Function(_typing___ushort2half_rd)) - - -@register -class _typing___ushort2half_ru(ConcreteTemplate): - key = globals()["__ushort2half_ru"] - cases = [signature(_type___half, uint16)] - - -register_global(__ushort2half_ru, types.Function(_typing___ushort2half_ru)) - - -@register -class _typing___half2ull_rn(ConcreteTemplate): - key = globals()["__half2ull_rn"] - cases = [signature(uint64, _type___half)] - - -register_global(__half2ull_rn, types.Function(_typing___half2ull_rn)) - - -@register -class _typing___half2ull_rd(ConcreteTemplate): - key = globals()["__half2ull_rd"] - cases = [signature(uint64, _type___half)] - - -register_global(__half2ull_rd, types.Function(_typing___half2ull_rd)) - - -@register -class _typing___half2ull_ru(ConcreteTemplate): - key = globals()["__half2ull_ru"] - cases = [signature(uint64, _type___half)] - - -register_global(__half2ull_ru, types.Function(_typing___half2ull_ru)) - - -@register -class _typing___ull2half_rn(ConcreteTemplate): - key = globals()["__ull2half_rn"] - cases = [signature(_type___half, uint64)] - - -register_global(__ull2half_rn, types.Function(_typing___ull2half_rn)) - - -@register -class _typing___ull2half_rz(ConcreteTemplate): - key = globals()["__ull2half_rz"] - cases = [signature(_type___half, uint64)] - - -register_global(__ull2half_rz, types.Function(_typing___ull2half_rz)) - - -@register -class _typing___ull2half_rd(ConcreteTemplate): - key = globals()["__ull2half_rd"] - cases = [signature(_type___half, uint64)] - - -register_global(__ull2half_rd, types.Function(_typing___ull2half_rd)) - - -@register -class _typing___ull2half_ru(ConcreteTemplate): - key = globals()["__ull2half_ru"] - cases = [signature(_type___half, uint64)] - - -register_global(__ull2half_ru, types.Function(_typing___ull2half_ru)) - - -@register -class _typing___half2ll_rn(ConcreteTemplate): - key = globals()["__half2ll_rn"] - cases = [signature(int64, _type___half)] - - -register_global(__half2ll_rn, types.Function(_typing___half2ll_rn)) - - -@register -class _typing___half2ll_rd(ConcreteTemplate): - key = globals()["__half2ll_rd"] - cases = [signature(int64, _type___half)] - - -register_global(__half2ll_rd, types.Function(_typing___half2ll_rd)) - - -@register -class _typing___half2ll_ru(ConcreteTemplate): - key = globals()["__half2ll_ru"] - cases = [signature(int64, _type___half)] - - -register_global(__half2ll_ru, types.Function(_typing___half2ll_ru)) - - -@register -class _typing___ll2half_rn(ConcreteTemplate): - key = globals()["__ll2half_rn"] - cases = [signature(_type___half, int64)] - - -register_global(__ll2half_rn, types.Function(_typing___ll2half_rn)) - - -@register -class _typing___ll2half_rz(ConcreteTemplate): - key = globals()["__ll2half_rz"] - cases = [signature(_type___half, int64)] - - -register_global(__ll2half_rz, types.Function(_typing___ll2half_rz)) - - -@register -class _typing___ll2half_rd(ConcreteTemplate): - key = globals()["__ll2half_rd"] - cases = [signature(_type___half, int64)] - - -register_global(__ll2half_rd, types.Function(_typing___ll2half_rd)) - - -@register -class _typing___ll2half_ru(ConcreteTemplate): - key = globals()["__ll2half_ru"] - cases = [signature(_type___half, int64)] - - -register_global(__ll2half_ru, types.Function(_typing___ll2half_ru)) - - -@register -class _typing_htrunc(ConcreteTemplate): - key = globals()["htrunc"] - cases = [signature(_type___half, _type___half)] - - -register_global(htrunc, types.Function(_typing_htrunc)) - - -@register -class _typing_hceil(ConcreteTemplate): - key = globals()["hceil"] - cases = [signature(_type___half, _type___half)] - - -register_global(hceil, types.Function(_typing_hceil)) - - -@register -class _typing_hfloor(ConcreteTemplate): - key = globals()["hfloor"] - cases = [signature(_type___half, _type___half)] - - -register_global(hfloor, types.Function(_typing_hfloor)) - - -@register -class _typing_hrint(ConcreteTemplate): - key = globals()["hrint"] - cases = [signature(_type___half, _type___half)] - - -register_global(hrint, types.Function(_typing_hrint)) - - -@register -class _typing_h2trunc(ConcreteTemplate): - key = globals()["h2trunc"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2trunc, types.Function(_typing_h2trunc)) - - -@register -class _typing_h2ceil(ConcreteTemplate): - key = globals()["h2ceil"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2ceil, types.Function(_typing_h2ceil)) - - -@register -class _typing_h2floor(ConcreteTemplate): - key = globals()["h2floor"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2floor, types.Function(_typing_h2floor)) - - -@register -class _typing_h2rint(ConcreteTemplate): - key = globals()["h2rint"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2rint, types.Function(_typing_h2rint)) - - -@register -class _typing___half2half2(ConcreteTemplate): - key = globals()["__half2half2"] - cases = [signature(_type___half2, _type___half)] - - -register_global(__half2half2, types.Function(_typing___half2half2)) - - -@register -class _typing___lowhigh2highlow(ConcreteTemplate): - key = globals()["__lowhigh2highlow"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(__lowhigh2highlow, types.Function(_typing___lowhigh2highlow)) - - -@register -class _typing___lows2half2(ConcreteTemplate): - key = globals()["__lows2half2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] - - -register_global(__lows2half2, types.Function(_typing___lows2half2)) - - -@register -class _typing___highs2half2(ConcreteTemplate): - key = globals()["__highs2half2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] - - -register_global(__highs2half2, types.Function(_typing___highs2half2)) - - -@register -class _typing___high2half(ConcreteTemplate): - key = globals()["__high2half"] - cases = [signature(_type___half, _type___half2)] - - -register_global(__high2half, types.Function(_typing___high2half)) - - -@register -class _typing___low2half(ConcreteTemplate): - key = globals()["__low2half"] - cases = [signature(_type___half, _type___half2)] - - -register_global(__low2half, types.Function(_typing___low2half)) +register_global(__half2int_rd, types.Function(_typing___half2int_rd)) @register -class _typing___hisinf(ConcreteTemplate): - key = globals()["__hisinf"] +class _typing___half2int_ru(ConcreteTemplate): + key = globals()["__half2int_ru"] cases = [signature(int32, _type___half)] -register_global(__hisinf, types.Function(_typing___hisinf)) +register_global(__half2int_ru, types.Function(_typing___half2int_ru)) @register -class _typing___halves2half2(ConcreteTemplate): - key = globals()["__halves2half2"] - cases = [signature(_type___half2, _type___half, _type___half)] +class _typing___int2half_rn(ConcreteTemplate): + key = globals()["__int2half_rn"] + cases = [signature(_type___half, int32)] -register_global(__halves2half2, types.Function(_typing___halves2half2)) +register_global(__int2half_rn, types.Function(_typing___int2half_rn)) @register -class _typing___low2half2(ConcreteTemplate): - key = globals()["__low2half2"] - cases = [signature(_type___half2, _type___half2)] +class _typing___int2half_rz(ConcreteTemplate): + key = globals()["__int2half_rz"] + cases = [signature(_type___half, int32)] -register_global(__low2half2, types.Function(_typing___low2half2)) +register_global(__int2half_rz, types.Function(_typing___int2half_rz)) @register -class _typing___high2half2(ConcreteTemplate): - key = globals()["__high2half2"] - cases = [signature(_type___half2, _type___half2)] +class _typing___int2half_rd(ConcreteTemplate): + key = globals()["__int2half_rd"] + cases = [signature(_type___half, int32)] -register_global(__high2half2, types.Function(_typing___high2half2)) +register_global(__int2half_rd, types.Function(_typing___int2half_rd)) @register -class _typing___half_as_short(ConcreteTemplate): - key = globals()["__half_as_short"] - cases = [signature(int16, _type___half)] +class _typing___int2half_ru(ConcreteTemplate): + key = globals()["__int2half_ru"] + cases = [signature(_type___half, int32)] -register_global(__half_as_short, types.Function(_typing___half_as_short)) +register_global(__int2half_ru, types.Function(_typing___int2half_ru)) @register -class _typing___half_as_ushort(ConcreteTemplate): - key = globals()["__half_as_ushort"] - cases = [signature(uint16, _type___half)] +class _typing___half2short_rn(ConcreteTemplate): + key = globals()["__half2short_rn"] + cases = [signature(int16, _type___half)] -register_global(__half_as_ushort, types.Function(_typing___half_as_ushort)) +register_global(__half2short_rn, types.Function(_typing___half2short_rn)) @register -class _typing___short_as_half(ConcreteTemplate): - key = globals()["__short_as_half"] - cases = [signature(_type___half, int16)] +class _typing___half2short_rd(ConcreteTemplate): + key = globals()["__half2short_rd"] + cases = [signature(int16, _type___half)] -register_global(__short_as_half, types.Function(_typing___short_as_half)) +register_global(__half2short_rd, types.Function(_typing___half2short_rd)) @register -class _typing___ushort_as_half(ConcreteTemplate): - key = globals()["__ushort_as_half"] - cases = [signature(_type___half, uint16)] +class _typing___half2short_ru(ConcreteTemplate): + key = globals()["__half2short_ru"] + cases = [signature(int16, _type___half)] -register_global(__ushort_as_half, types.Function(_typing___ushort_as_half)) +register_global(__half2short_ru, types.Function(_typing___half2short_ru)) @register -class _typing___hmax(ConcreteTemplate): - key = globals()["__hmax"] - cases = [signature(_type___half, _type___half, _type___half)] +class _typing___short2half_rn(ConcreteTemplate): + key = globals()["__short2half_rn"] + cases = [signature(_type___half, int16)] -register_global(__hmax, types.Function(_typing___hmax)) +register_global(__short2half_rn, types.Function(_typing___short2half_rn)) @register -class _typing___hmin(ConcreteTemplate): - key = globals()["__hmin"] - cases = [signature(_type___half, _type___half, _type___half)] +class _typing___short2half_rz(ConcreteTemplate): + key = globals()["__short2half_rz"] + cases = [signature(_type___half, int16)] -register_global(__hmin, types.Function(_typing___hmin)) +register_global(__short2half_rz, types.Function(_typing___short2half_rz)) @register -class _typing___hmax2(ConcreteTemplate): - key = globals()["__hmax2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___short2half_rd(ConcreteTemplate): + key = globals()["__short2half_rd"] + cases = [signature(_type___half, int16)] -register_global(__hmax2, types.Function(_typing___hmax2)) +register_global(__short2half_rd, types.Function(_typing___short2half_rd)) @register -class _typing___hmin2(ConcreteTemplate): - key = globals()["__hmin2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___short2half_ru(ConcreteTemplate): + key = globals()["__short2half_ru"] + cases = [signature(_type___half, int16)] -register_global(__hmin2, types.Function(_typing___hmin2)) +register_global(__short2half_ru, types.Function(_typing___short2half_ru)) @register -class _typing___shfl_sync(ConcreteTemplate): - key = globals()["__shfl_sync"] - cases = [ - signature(_type___half2, uint32, _type___half2, int32, int32), - signature(_type___half, uint32, _type___half, int32, int32), - ] +class _typing___half2uint_rn(ConcreteTemplate): + key = globals()["__half2uint_rn"] + cases = [signature(uint32, _type___half)] -register_global(__shfl_sync, types.Function(_typing___shfl_sync)) +register_global(__half2uint_rn, types.Function(_typing___half2uint_rn)) @register -class _typing___shfl_up_sync(ConcreteTemplate): - key = globals()["__shfl_up_sync"] - cases = [ - signature(_type___half2, uint32, _type___half2, uint32, int32), - signature(_type___half, uint32, _type___half, uint32, int32), - ] +class _typing___half2uint_rd(ConcreteTemplate): + key = globals()["__half2uint_rd"] + cases = [signature(uint32, _type___half)] -register_global(__shfl_up_sync, types.Function(_typing___shfl_up_sync)) +register_global(__half2uint_rd, types.Function(_typing___half2uint_rd)) @register -class _typing___shfl_down_sync(ConcreteTemplate): - key = globals()["__shfl_down_sync"] - cases = [ - signature(_type___half2, uint32, _type___half2, uint32, int32), - signature(_type___half, uint32, _type___half, uint32, int32), - ] +class _typing___half2uint_ru(ConcreteTemplate): + key = globals()["__half2uint_ru"] + cases = [signature(uint32, _type___half)] -register_global(__shfl_down_sync, types.Function(_typing___shfl_down_sync)) +register_global(__half2uint_ru, types.Function(_typing___half2uint_ru)) @register -class _typing___shfl_xor_sync(ConcreteTemplate): - key = globals()["__shfl_xor_sync"] - cases = [ - signature(_type___half2, uint32, _type___half2, int32, int32), - signature(_type___half, uint32, _type___half, int32, int32), - ] +class _typing___uint2half_rn(ConcreteTemplate): + key = globals()["__uint2half_rn"] + cases = [signature(_type___half, uint32)] -register_global(__shfl_xor_sync, types.Function(_typing___shfl_xor_sync)) +register_global(__uint2half_rn, types.Function(_typing___uint2half_rn)) @register -class _typing___ldg(ConcreteTemplate): - key = globals()["__ldg"] - cases = [ - signature(_type___half2, CPointer(_type___half2)), - signature(_type___half, CPointer(_type___half)), - ] +class _typing___uint2half_rz(ConcreteTemplate): + key = globals()["__uint2half_rz"] + cases = [signature(_type___half, uint32)] -register_global(__ldg, types.Function(_typing___ldg)) +register_global(__uint2half_rz, types.Function(_typing___uint2half_rz)) @register -class _typing___ldcg(ConcreteTemplate): - key = globals()["__ldcg"] - cases = [ - signature(_type___half2, CPointer(_type___half2)), - signature(_type___half, CPointer(_type___half)), - ] +class _typing___uint2half_rd(ConcreteTemplate): + key = globals()["__uint2half_rd"] + cases = [signature(_type___half, uint32)] -register_global(__ldcg, types.Function(_typing___ldcg)) +register_global(__uint2half_rd, types.Function(_typing___uint2half_rd)) @register -class _typing___ldca(ConcreteTemplate): - key = globals()["__ldca"] - cases = [ - signature(_type___half2, CPointer(_type___half2)), - signature(_type___half, CPointer(_type___half)), - ] +class _typing___uint2half_ru(ConcreteTemplate): + key = globals()["__uint2half_ru"] + cases = [signature(_type___half, uint32)] -register_global(__ldca, types.Function(_typing___ldca)) +register_global(__uint2half_ru, types.Function(_typing___uint2half_ru)) @register -class _typing___ldcs(ConcreteTemplate): - key = globals()["__ldcs"] - cases = [ - signature(_type___half2, CPointer(_type___half2)), - signature(_type___half, CPointer(_type___half)), - ] +class _typing___half2ushort_rn(ConcreteTemplate): + key = globals()["__half2ushort_rn"] + cases = [signature(uint16, _type___half)] -register_global(__ldcs, types.Function(_typing___ldcs)) +register_global(__half2ushort_rn, types.Function(_typing___half2ushort_rn)) @register -class _typing___ldlu(ConcreteTemplate): - key = globals()["__ldlu"] - cases = [ - signature(_type___half2, CPointer(_type___half2)), - signature(_type___half, CPointer(_type___half)), - ] +class _typing___half2ushort_rd(ConcreteTemplate): + key = globals()["__half2ushort_rd"] + cases = [signature(uint16, _type___half)] -register_global(__ldlu, types.Function(_typing___ldlu)) +register_global(__half2ushort_rd, types.Function(_typing___half2ushort_rd)) @register -class _typing___ldcv(ConcreteTemplate): - key = globals()["__ldcv"] - cases = [ - signature(_type___half2, CPointer(_type___half2)), - signature(_type___half, CPointer(_type___half)), - ] +class _typing___half2ushort_ru(ConcreteTemplate): + key = globals()["__half2ushort_ru"] + cases = [signature(uint16, _type___half)] -register_global(__ldcv, types.Function(_typing___ldcv)) +register_global(__half2ushort_ru, types.Function(_typing___half2ushort_ru)) @register -class _typing___stwb(ConcreteTemplate): - key = globals()["__stwb"] - cases = [ - signature(void, CPointer(_type___half2), _type___half2), - signature(void, CPointer(_type___half), _type___half), - ] +class _typing___ushort2half_rn(ConcreteTemplate): + key = globals()["__ushort2half_rn"] + cases = [signature(_type___half, uint16)] -register_global(__stwb, types.Function(_typing___stwb)) +register_global(__ushort2half_rn, types.Function(_typing___ushort2half_rn)) @register -class _typing___stcg(ConcreteTemplate): - key = globals()["__stcg"] - cases = [ - signature(void, CPointer(_type___half2), _type___half2), - signature(void, CPointer(_type___half), _type___half), - ] +class _typing___ushort2half_rz(ConcreteTemplate): + key = globals()["__ushort2half_rz"] + cases = [signature(_type___half, uint16)] -register_global(__stcg, types.Function(_typing___stcg)) +register_global(__ushort2half_rz, types.Function(_typing___ushort2half_rz)) @register -class _typing___stcs(ConcreteTemplate): - key = globals()["__stcs"] - cases = [ - signature(void, CPointer(_type___half2), _type___half2), - signature(void, CPointer(_type___half), _type___half), - ] +class _typing___ushort2half_rd(ConcreteTemplate): + key = globals()["__ushort2half_rd"] + cases = [signature(_type___half, uint16)] -register_global(__stcs, types.Function(_typing___stcs)) +register_global(__ushort2half_rd, types.Function(_typing___ushort2half_rd)) @register -class _typing___stwt(ConcreteTemplate): - key = globals()["__stwt"] - cases = [ - signature(void, CPointer(_type___half2), _type___half2), - signature(void, CPointer(_type___half), _type___half), - ] +class _typing___ushort2half_ru(ConcreteTemplate): + key = globals()["__ushort2half_ru"] + cases = [signature(_type___half, uint16)] -register_global(__stwt, types.Function(_typing___stwt)) +register_global(__ushort2half_ru, types.Function(_typing___ushort2half_ru)) @register -class _typing___heq2(ConcreteTemplate): - key = globals()["__heq2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___half2ull_rn(ConcreteTemplate): + key = globals()["__half2ull_rn"] + cases = [signature(uint64, _type___half)] -register_global(__heq2, types.Function(_typing___heq2)) +register_global(__half2ull_rn, types.Function(_typing___half2ull_rn)) @register -class _typing___hne2(ConcreteTemplate): - key = globals()["__hne2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___half2ull_rd(ConcreteTemplate): + key = globals()["__half2ull_rd"] + cases = [signature(uint64, _type___half)] -register_global(__hne2, types.Function(_typing___hne2)) +register_global(__half2ull_rd, types.Function(_typing___half2ull_rd)) @register -class _typing___hle2(ConcreteTemplate): - key = globals()["__hle2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___half2ull_ru(ConcreteTemplate): + key = globals()["__half2ull_ru"] + cases = [signature(uint64, _type___half)] -register_global(__hle2, types.Function(_typing___hle2)) +register_global(__half2ull_ru, types.Function(_typing___half2ull_ru)) @register -class _typing___hge2(ConcreteTemplate): - key = globals()["__hge2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ull2half_rn(ConcreteTemplate): + key = globals()["__ull2half_rn"] + cases = [signature(_type___half, uint64)] -register_global(__hge2, types.Function(_typing___hge2)) +register_global(__ull2half_rn, types.Function(_typing___ull2half_rn)) @register -class _typing___hlt2(ConcreteTemplate): - key = globals()["__hlt2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ull2half_rz(ConcreteTemplate): + key = globals()["__ull2half_rz"] + cases = [signature(_type___half, uint64)] -register_global(__hlt2, types.Function(_typing___hlt2)) +register_global(__ull2half_rz, types.Function(_typing___ull2half_rz)) @register -class _typing___hgt2(ConcreteTemplate): - key = globals()["__hgt2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ull2half_rd(ConcreteTemplate): + key = globals()["__ull2half_rd"] + cases = [signature(_type___half, uint64)] -register_global(__hgt2, types.Function(_typing___hgt2)) +register_global(__ull2half_rd, types.Function(_typing___ull2half_rd)) @register -class _typing___hequ2(ConcreteTemplate): - key = globals()["__hequ2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ull2half_ru(ConcreteTemplate): + key = globals()["__ull2half_ru"] + cases = [signature(_type___half, uint64)] -register_global(__hequ2, types.Function(_typing___hequ2)) +register_global(__ull2half_ru, types.Function(_typing___ull2half_ru)) @register -class _typing___hneu2(ConcreteTemplate): - key = globals()["__hneu2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___half2ll_rn(ConcreteTemplate): + key = globals()["__half2ll_rn"] + cases = [signature(int64, _type___half)] -register_global(__hneu2, types.Function(_typing___hneu2)) +register_global(__half2ll_rn, types.Function(_typing___half2ll_rn)) @register -class _typing___hleu2(ConcreteTemplate): - key = globals()["__hleu2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___half2ll_rd(ConcreteTemplate): + key = globals()["__half2ll_rd"] + cases = [signature(int64, _type___half)] -register_global(__hleu2, types.Function(_typing___hleu2)) +register_global(__half2ll_rd, types.Function(_typing___half2ll_rd)) @register -class _typing___hgeu2(ConcreteTemplate): - key = globals()["__hgeu2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___half2ll_ru(ConcreteTemplate): + key = globals()["__half2ll_ru"] + cases = [signature(int64, _type___half)] -register_global(__hgeu2, types.Function(_typing___hgeu2)) +register_global(__half2ll_ru, types.Function(_typing___half2ll_ru)) @register -class _typing___hltu2(ConcreteTemplate): - key = globals()["__hltu2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ll2half_rn(ConcreteTemplate): + key = globals()["__ll2half_rn"] + cases = [signature(_type___half, int64)] -register_global(__hltu2, types.Function(_typing___hltu2)) +register_global(__ll2half_rn, types.Function(_typing___ll2half_rn)) @register -class _typing___hgtu2(ConcreteTemplate): - key = globals()["__hgtu2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ll2half_rz(ConcreteTemplate): + key = globals()["__ll2half_rz"] + cases = [signature(_type___half, int64)] -register_global(__hgtu2, types.Function(_typing___hgtu2)) +register_global(__ll2half_rz, types.Function(_typing___ll2half_rz)) @register -class _typing___heq2_mask(ConcreteTemplate): - key = globals()["__heq2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___ll2half_rd(ConcreteTemplate): + key = globals()["__ll2half_rd"] + cases = [signature(_type___half, int64)] -register_global(__heq2_mask, types.Function(_typing___heq2_mask)) +register_global(__ll2half_rd, types.Function(_typing___ll2half_rd)) @register -class _typing___hne2_mask(ConcreteTemplate): - key = globals()["__hne2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___ll2half_ru(ConcreteTemplate): + key = globals()["__ll2half_ru"] + cases = [signature(_type___half, int64)] -register_global(__hne2_mask, types.Function(_typing___hne2_mask)) +register_global(__ll2half_ru, types.Function(_typing___ll2half_ru)) @register -class _typing___hle2_mask(ConcreteTemplate): - key = globals()["__hle2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing_htrunc(ConcreteTemplate): + key = globals()["htrunc"] + cases = [signature(_type___half, _type___half)] -register_global(__hle2_mask, types.Function(_typing___hle2_mask)) +register_global(htrunc, types.Function(_typing_htrunc)) @register -class _typing___hge2_mask(ConcreteTemplate): - key = globals()["__hge2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing_hceil(ConcreteTemplate): + key = globals()["hceil"] + cases = [signature(_type___half, _type___half)] -register_global(__hge2_mask, types.Function(_typing___hge2_mask)) +register_global(hceil, types.Function(_typing_hceil)) @register -class _typing___hlt2_mask(ConcreteTemplate): - key = globals()["__hlt2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing_hfloor(ConcreteTemplate): + key = globals()["hfloor"] + cases = [signature(_type___half, _type___half)] -register_global(__hlt2_mask, types.Function(_typing___hlt2_mask)) +register_global(hfloor, types.Function(_typing_hfloor)) @register -class _typing___hgt2_mask(ConcreteTemplate): - key = globals()["__hgt2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing_hrint(ConcreteTemplate): + key = globals()["hrint"] + cases = [signature(_type___half, _type___half)] -register_global(__hgt2_mask, types.Function(_typing___hgt2_mask)) +register_global(hrint, types.Function(_typing_hrint)) @register -class _typing___hequ2_mask(ConcreteTemplate): - key = globals()["__hequ2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___hisinf(ConcreteTemplate): + key = globals()["__hisinf"] + cases = [signature(int32, _type___half)] -register_global(__hequ2_mask, types.Function(_typing___hequ2_mask)) +register_global(__hisinf, types.Function(_typing___hisinf)) @register -class _typing___hneu2_mask(ConcreteTemplate): - key = globals()["__hneu2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___half_as_short(ConcreteTemplate): + key = globals()["__half_as_short"] + cases = [signature(int16, _type___half)] -register_global(__hneu2_mask, types.Function(_typing___hneu2_mask)) +register_global(__half_as_short, types.Function(_typing___half_as_short)) @register -class _typing___hleu2_mask(ConcreteTemplate): - key = globals()["__hleu2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___half_as_ushort(ConcreteTemplate): + key = globals()["__half_as_ushort"] + cases = [signature(uint16, _type___half)] -register_global(__hleu2_mask, types.Function(_typing___hleu2_mask)) +register_global(__half_as_ushort, types.Function(_typing___half_as_ushort)) @register -class _typing___hgeu2_mask(ConcreteTemplate): - key = globals()["__hgeu2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___short_as_half(ConcreteTemplate): + key = globals()["__short_as_half"] + cases = [signature(_type___half, int16)] -register_global(__hgeu2_mask, types.Function(_typing___hgeu2_mask)) +register_global(__short_as_half, types.Function(_typing___short_as_half)) @register -class _typing___hltu2_mask(ConcreteTemplate): - key = globals()["__hltu2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___ushort_as_half(ConcreteTemplate): + key = globals()["__ushort_as_half"] + cases = [signature(_type___half, uint16)] -register_global(__hltu2_mask, types.Function(_typing___hltu2_mask)) +register_global(__ushort_as_half, types.Function(_typing___ushort_as_half)) @register -class _typing___hgtu2_mask(ConcreteTemplate): - key = globals()["__hgtu2_mask"] - cases = [signature(uint32, _type___half2, _type___half2)] +class _typing___hmax(ConcreteTemplate): + key = globals()["__hmax"] + cases = [signature(_type___half, _type___half, _type___half)] -register_global(__hgtu2_mask, types.Function(_typing___hgtu2_mask)) +register_global(__hmax, types.Function(_typing___hmax)) @register -class _typing___hisnan2(ConcreteTemplate): - key = globals()["__hisnan2"] - cases = [signature(_type___half2, _type___half2)] +class _typing___hmin(ConcreteTemplate): + key = globals()["__hmin"] + cases = [signature(_type___half, _type___half, _type___half)] -register_global(__hisnan2, types.Function(_typing___hisnan2)) +register_global(__hmin, types.Function(_typing___hmin)) @register -class _typing___hadd2(ConcreteTemplate): - key = globals()["__hadd2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___shfl_sync(ConcreteTemplate): + key = globals()["__shfl_sync"] + cases = [signature(_type___half, uint32, _type___half, int32, int32)] -register_global(__hadd2, types.Function(_typing___hadd2)) +register_global(__shfl_sync, types.Function(_typing___shfl_sync)) @register -class _typing___hsub2(ConcreteTemplate): - key = globals()["__hsub2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___shfl_up_sync(ConcreteTemplate): + key = globals()["__shfl_up_sync"] + cases = [signature(_type___half, uint32, _type___half, uint32, int32)] -register_global(__hsub2, types.Function(_typing___hsub2)) +register_global(__shfl_up_sync, types.Function(_typing___shfl_up_sync)) @register -class _typing___hmul2(ConcreteTemplate): - key = globals()["__hmul2"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___shfl_down_sync(ConcreteTemplate): + key = globals()["__shfl_down_sync"] + cases = [signature(_type___half, uint32, _type___half, uint32, int32)] -register_global(__hmul2, types.Function(_typing___hmul2)) +register_global(__shfl_down_sync, types.Function(_typing___shfl_down_sync)) @register -class _typing___hadd2_rn(ConcreteTemplate): - key = globals()["__hadd2_rn"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___shfl_xor_sync(ConcreteTemplate): + key = globals()["__shfl_xor_sync"] + cases = [signature(_type___half, uint32, _type___half, int32, int32)] -register_global(__hadd2_rn, types.Function(_typing___hadd2_rn)) +register_global(__shfl_xor_sync, types.Function(_typing___shfl_xor_sync)) @register -class _typing___hsub2_rn(ConcreteTemplate): - key = globals()["__hsub2_rn"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ldg(ConcreteTemplate): + key = globals()["__ldg"] + cases = [signature(_type___half, CPointer(_type___half))] -register_global(__hsub2_rn, types.Function(_typing___hsub2_rn)) +register_global(__ldg, types.Function(_typing___ldg)) @register -class _typing___hmul2_rn(ConcreteTemplate): - key = globals()["__hmul2_rn"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ldcg(ConcreteTemplate): + key = globals()["__ldcg"] + cases = [signature(_type___half, CPointer(_type___half))] -register_global(__hmul2_rn, types.Function(_typing___hmul2_rn)) +register_global(__ldcg, types.Function(_typing___ldcg)) @register -class _typing___h2div(ConcreteTemplate): - key = globals()["__h2div"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ldca(ConcreteTemplate): + key = globals()["__ldca"] + cases = [signature(_type___half, CPointer(_type___half))] -register_global(__h2div, types.Function(_typing___h2div)) +register_global(__ldca, types.Function(_typing___ldca)) @register -class _typing___habs2(ConcreteTemplate): - key = globals()["__habs2"] - cases = [signature(_type___half2, _type___half2)] +class _typing___ldcs(ConcreteTemplate): + key = globals()["__ldcs"] + cases = [signature(_type___half, CPointer(_type___half))] -register_global(__habs2, types.Function(_typing___habs2)) +register_global(__ldcs, types.Function(_typing___ldcs)) @register -class _typing___hadd2_sat(ConcreteTemplate): - key = globals()["__hadd2_sat"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ldlu(ConcreteTemplate): + key = globals()["__ldlu"] + cases = [signature(_type___half, CPointer(_type___half))] -register_global(__hadd2_sat, types.Function(_typing___hadd2_sat)) +register_global(__ldlu, types.Function(_typing___ldlu)) @register -class _typing___hsub2_sat(ConcreteTemplate): - key = globals()["__hsub2_sat"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___ldcv(ConcreteTemplate): + key = globals()["__ldcv"] + cases = [signature(_type___half, CPointer(_type___half))] -register_global(__hsub2_sat, types.Function(_typing___hsub2_sat)) +register_global(__ldcv, types.Function(_typing___ldcv)) @register -class _typing___hmul2_sat(ConcreteTemplate): - key = globals()["__hmul2_sat"] - cases = [signature(_type___half2, _type___half2, _type___half2)] +class _typing___stwb(ConcreteTemplate): + key = globals()["__stwb"] + cases = [signature(void, CPointer(_type___half), _type___half)] -register_global(__hmul2_sat, types.Function(_typing___hmul2_sat)) +register_global(__stwb, types.Function(_typing___stwb)) @register -class _typing___hfma2(ConcreteTemplate): - key = globals()["__hfma2"] - cases = [ - signature(_type___half2, _type___half2, _type___half2, _type___half2) - ] +class _typing___stcg(ConcreteTemplate): + key = globals()["__stcg"] + cases = [signature(void, CPointer(_type___half), _type___half)] -register_global(__hfma2, types.Function(_typing___hfma2)) +register_global(__stcg, types.Function(_typing___stcg)) @register -class _typing___hfma2_sat(ConcreteTemplate): - key = globals()["__hfma2_sat"] - cases = [ - signature(_type___half2, _type___half2, _type___half2, _type___half2) - ] +class _typing___stcs(ConcreteTemplate): + key = globals()["__stcs"] + cases = [signature(void, CPointer(_type___half), _type___half)] -register_global(__hfma2_sat, types.Function(_typing___hfma2_sat)) +register_global(__stcs, types.Function(_typing___stcs)) @register -class _typing___hneg2(ConcreteTemplate): - key = globals()["__hneg2"] - cases = [signature(_type___half2, _type___half2)] +class _typing___stwt(ConcreteTemplate): + key = globals()["__stwt"] + cases = [signature(void, CPointer(_type___half), _type___half)] -register_global(__hneg2, types.Function(_typing___hneg2)) +register_global(__stwt, types.Function(_typing___stwt)) @register @@ -13248,114 +7525,6 @@ class _typing___hneg(ConcreteTemplate): register_global(__hneg, types.Function(_typing___hneg)) -@register -class _typing___hbeq2(ConcreteTemplate): - key = globals()["__hbeq2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbeq2, types.Function(_typing___hbeq2)) - - -@register -class _typing___hbne2(ConcreteTemplate): - key = globals()["__hbne2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbne2, types.Function(_typing___hbne2)) - - -@register -class _typing___hble2(ConcreteTemplate): - key = globals()["__hble2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hble2, types.Function(_typing___hble2)) - - -@register -class _typing___hbge2(ConcreteTemplate): - key = globals()["__hbge2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbge2, types.Function(_typing___hbge2)) - - -@register -class _typing___hblt2(ConcreteTemplate): - key = globals()["__hblt2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hblt2, types.Function(_typing___hblt2)) - - -@register -class _typing___hbgt2(ConcreteTemplate): - key = globals()["__hbgt2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbgt2, types.Function(_typing___hbgt2)) - - -@register -class _typing___hbequ2(ConcreteTemplate): - key = globals()["__hbequ2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbequ2, types.Function(_typing___hbequ2)) - - -@register -class _typing___hbneu2(ConcreteTemplate): - key = globals()["__hbneu2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbneu2, types.Function(_typing___hbneu2)) - - -@register -class _typing___hbleu2(ConcreteTemplate): - key = globals()["__hbleu2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbleu2, types.Function(_typing___hbleu2)) - - -@register -class _typing___hbgeu2(ConcreteTemplate): - key = globals()["__hbgeu2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbgeu2, types.Function(_typing___hbgeu2)) - - -@register -class _typing___hbltu2(ConcreteTemplate): - key = globals()["__hbltu2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbltu2, types.Function(_typing___hbltu2)) - - -@register -class _typing___hbgtu2(ConcreteTemplate): - key = globals()["__hbgtu2"] - cases = [signature(bool_, _type___half2, _type___half2)] - - -register_global(__hbgtu2, types.Function(_typing___hbgtu2)) - - @register class _typing___heq(ConcreteTemplate): key = globals()["__heq"] @@ -13500,46 +7669,6 @@ class _typing___hfma_relu(ConcreteTemplate): register_global(__hfma_relu, types.Function(_typing___hfma_relu)) -@register -class _typing___hmax2_nan(ConcreteTemplate): - key = globals()["__hmax2_nan"] - cases = [signature(_type___half2, _type___half2, _type___half2)] - - -register_global(__hmax2_nan, types.Function(_typing___hmax2_nan)) - - -@register -class _typing___hmin2_nan(ConcreteTemplate): - key = globals()["__hmin2_nan"] - cases = [signature(_type___half2, _type___half2, _type___half2)] - - -register_global(__hmin2_nan, types.Function(_typing___hmin2_nan)) - - -@register -class _typing___hfma2_relu(ConcreteTemplate): - key = globals()["__hfma2_relu"] - cases = [ - signature(_type___half2, _type___half2, _type___half2, _type___half2) - ] - - -register_global(__hfma2_relu, types.Function(_typing___hfma2_relu)) - - -@register -class _typing___hcmadd(ConcreteTemplate): - key = globals()["__hcmadd"] - cases = [ - signature(_type___half2, _type___half2, _type___half2, _type___half2) - ] - - -register_global(__hcmadd, types.Function(_typing___hcmadd)) - - @register class _typing_hsqrt(ConcreteTemplate): key = globals()["hsqrt"] @@ -13612,15 +7741,6 @@ class _typing_htanh_approx(ConcreteTemplate): register_global(htanh_approx, types.Function(_typing_htanh_approx)) -@register -class _typing_h2tanh_approx(ConcreteTemplate): - key = globals()["h2tanh_approx"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2tanh_approx, types.Function(_typing_h2tanh_approx)) - - @register class _typing_htanh(ConcreteTemplate): key = globals()["htanh"] @@ -13630,15 +7750,6 @@ class _typing_htanh(ConcreteTemplate): register_global(htanh, types.Function(_typing_htanh)) -@register -class _typing_h2tanh(ConcreteTemplate): - key = globals()["h2tanh"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2tanh, types.Function(_typing_h2tanh)) - - @register class _typing_hexp2(ConcreteTemplate): key = globals()["hexp2"] @@ -13675,112 +7786,10 @@ class _typing_hsin(ConcreteTemplate): register_global(hsin, types.Function(_typing_hsin)) -@register -class _typing_h2sqrt(ConcreteTemplate): - key = globals()["h2sqrt"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2sqrt, types.Function(_typing_h2sqrt)) - - -@register -class _typing_h2rsqrt(ConcreteTemplate): - key = globals()["h2rsqrt"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2rsqrt, types.Function(_typing_h2rsqrt)) - - -@register -class _typing_h2rcp(ConcreteTemplate): - key = globals()["h2rcp"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2rcp, types.Function(_typing_h2rcp)) - - -@register -class _typing_h2log(ConcreteTemplate): - key = globals()["h2log"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2log, types.Function(_typing_h2log)) - - -@register -class _typing_h2log2(ConcreteTemplate): - key = globals()["h2log2"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2log2, types.Function(_typing_h2log2)) - - -@register -class _typing_h2log10(ConcreteTemplate): - key = globals()["h2log10"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2log10, types.Function(_typing_h2log10)) - - -@register -class _typing_h2exp(ConcreteTemplate): - key = globals()["h2exp"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2exp, types.Function(_typing_h2exp)) - - -@register -class _typing_h2exp2(ConcreteTemplate): - key = globals()["h2exp2"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2exp2, types.Function(_typing_h2exp2)) - - -@register -class _typing_h2exp10(ConcreteTemplate): - key = globals()["h2exp10"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2exp10, types.Function(_typing_h2exp10)) - - -@register -class _typing_h2cos(ConcreteTemplate): - key = globals()["h2cos"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2cos, types.Function(_typing_h2cos)) - - -@register -class _typing_h2sin(ConcreteTemplate): - key = globals()["h2sin"] - cases = [signature(_type___half2, _type___half2)] - - -register_global(h2sin, types.Function(_typing_h2sin)) - - @register class _typing_atomicAdd(ConcreteTemplate): key = globals()["atomicAdd"] - cases = [ - signature(_type___half2, CPointer(_type___half2), _type___half2), - signature(_type___half, CPointer(_type___half), _type___half), - ] + cases = [signature(_type___half, CPointer(_type___half), _type___half)] register_global(atomicAdd, types.Function(_typing_atomicAdd)) @@ -13932,30 +7941,26 @@ def _genfp16_binary_operator(op): # Aliases: -__half_raw = unnamed1302257 -__nv_half_raw = unnamed1302257 -__half2_raw = unnamed1302366 -__nv_half2_raw = unnamed1302366 +__half_raw = unnamed1362071 +__nv_half_raw = unnamed1362071 +__half2_raw = unnamed1362180 +__nv_half2_raw = unnamed1362180 half = __half __nv_half = __half nv_half = __half -half2 = __half2 -__nv_half2 = __half2 -nv_half2 = __half2 # Symbols: _NBTYPE_SYMBOLS = [ - "_type_unnamed1302257", - "_type_unnamed1302366", + "_type_unnamed1362071", + "_type_unnamed1362180", "_type___half", - "_type___half2", ] -_RECORD_SYMBOLS = ["unnamed1302257", "unnamed1302366", "__half", "__half2"] +_RECORD_SYMBOLS = ["unnamed1362071", "unnamed1362180", "__half"] _FUNCTION_SYMBOLS = [ @@ -13966,10 +7971,6 @@ def _genfp16_binary_operator(op): "__float2half_rd", "__float2half_ru", "__half2float", - "__float2half2_rn", - "__floats2half2_rn", - "__low2float", - "__high2float", "__half2char_rz", "__half2uchar_rz", "__half2short_rz", @@ -13978,9 +7979,6 @@ def _genfp16_binary_operator(op): "__half2uint_rz", "__half2ll_rz", "__half2ull_rz", - "make_half2", - "__float22half2_rn", - "__half22float2", "__half2int_rn", "__half2int_rd", "__half2int_ru", @@ -14027,95 +8025,27 @@ def _genfp16_binary_operator(op): "hceil", "hfloor", "hrint", - "h2trunc", - "h2ceil", - "h2floor", - "h2rint", - "__half2half2", - "__lowhigh2highlow", - "__lows2half2", - "__highs2half2", - "__high2half", - "__low2half", "__hisinf", - "__halves2half2", - "__low2half2", - "__high2half2", "__half_as_short", "__half_as_ushort", "__short_as_half", "__ushort_as_half", "__hmax", "__hmin", - "__hmax2", - "__hmin2", - "__shfl_sync", - "__shfl_up_sync", - "__shfl_down_sync", - "__shfl_xor_sync", "__shfl_sync", "__shfl_up_sync", "__shfl_down_sync", "__shfl_xor_sync", "__ldg", - "__ldg", "__ldcg", - "__ldcg", - "__ldca", "__ldca", "__ldcs", - "__ldcs", - "__ldlu", "__ldlu", "__ldcv", - "__ldcv", - "__stwb", "__stwb", "__stcg", - "__stcg", - "__stcs", "__stcs", "__stwt", - "__stwt", - "__heq2", - "__hne2", - "__hle2", - "__hge2", - "__hlt2", - "__hgt2", - "__hequ2", - "__hneu2", - "__hleu2", - "__hgeu2", - "__hltu2", - "__hgtu2", - "__heq2_mask", - "__hne2_mask", - "__hle2_mask", - "__hge2_mask", - "__hlt2_mask", - "__hgt2_mask", - "__hequ2_mask", - "__hneu2_mask", - "__hleu2_mask", - "__hgeu2_mask", - "__hltu2_mask", - "__hgtu2_mask", - "__hisnan2", - "__hadd2", - "__hsub2", - "__hmul2", - "__hadd2_rn", - "__hsub2_rn", - "__hmul2_rn", - "__h2div", - "__habs2", - "__hadd2_sat", - "__hsub2_sat", - "__hmul2_sat", - "__hfma2", - "__hfma2_sat", - "__hneg2", "__habs", "__hadd", "__hsub", @@ -14130,18 +8060,6 @@ def _genfp16_binary_operator(op): "__hfma", "__hfma_sat", "__hneg", - "__hbeq2", - "__hbne2", - "__hble2", - "__hbge2", - "__hblt2", - "__hbgt2", - "__hbequ2", - "__hbneu2", - "__hbleu2", - "__hbgeu2", - "__hbltu2", - "__hbgtu2", "__heq", "__hne", "__hle", @@ -14158,10 +8076,6 @@ def _genfp16_binary_operator(op): "__hmax_nan", "__hmin_nan", "__hfma_relu", - "__hmax2_nan", - "__hmin2_nan", - "__hfma2_relu", - "__hcmadd", "hsqrt", "hrsqrt", "hrcp", @@ -14170,25 +8084,11 @@ def _genfp16_binary_operator(op): "hlog10", "hexp", "htanh_approx", - "h2tanh_approx", "htanh", - "h2tanh", "hexp2", "hexp10", "hcos", "hsin", - "h2sqrt", - "h2rsqrt", - "h2rcp", - "h2log", - "h2log2", - "h2log10", - "h2exp", - "h2exp2", - "h2exp10", - "h2cos", - "h2sin", - "atomicAdd", "atomicAdd", "operator+", "operator-", @@ -14206,22 +8106,6 @@ def _genfp16_binary_operator(op): "operator<", "operator>=", "operator<=", - "operator+", - "operator-", - "operator*", - "operator/", - "operator+=", - "operator-=", - "operator*=", - "operator/=", - "operator+", - "operator-", - "operator==", - "operator!=", - "operator>", - "operator<", - "operator>=", - "operator<=", ] diff --git a/numba_cuda/numba/cuda/fp16.py b/numba_cuda/numba/cuda/fp16.py index 966f7fb22..ed9d73b60 100644 --- a/numba_cuda/numba/cuda/fp16.py +++ b/numba_cuda/numba/cuda/fp16.py @@ -73,7 +73,6 @@ __hadd, __hadd_rn, __hadd_sat, - __hcmadd, __hdiv as hdiv, __hdiv, __heq as heq, @@ -287,7 +286,6 @@ def exp2_ol(a): "__hadd", "__hadd_rn", "__hadd_sat", - "__hcmadd", "hdiv", "__hdiv", "heq",