[BUG] FP16 Type Fails with External Functions

[isVoid]

Describe the bug
If we use fp16 to foreign functions, there's prototype mismatch error. The error reproduces under:

1. ctk 12.8 ctypes linker
2. An older ctk with nvjitlink such as 12.0.1, 12.2, 11.8

Error Log:

/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/dispatcher.py:686: NumbaPerformanceWarning: Grid size 1 will likely result in GPU under-utilization due to low occupancy.
  warn(NumbaPerformanceWarning(msg))
Traceback (most recent call last):
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/driver.py", line 3247, in _add_data
    driver.cuLinkAddData(
    ~~~~~~~~~~~~~~~~~~~~^
        self.handle,
        ^^^^^^^^^^^^
    ...<6 lines>...
        None,
        ^^^^^
    )
    ^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/driver.py", line 374, in safe_cuda_api_call
    self._check_ctypes_error(fname, retcode)
    ~~~~~~~~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/driver.py", line 445, in _check_ctypes_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [300] Call to cuLinkAddData results in CUDA_ERROR_INVALID_SOURCE

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/wangm/scratch/072125/minrepro.py", line 29, in <module>
    fails[1, 1](arr)
    ~~~~~~~~~~~^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/dispatcher.py", line 689, in __call__
    return self.dispatcher.call(
           ~~~~~~~~~~~~~~~~~~~~^
        args, self.griddim, self.blockdim, self.stream, self.sharedmem
        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    )
    ^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/dispatcher.py", line 835, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/dispatcher.py", line 843, in _compile_for_args
    return self.compile(tuple(argtypes))
           ~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/dispatcher.py", line 1111, in compile
    kernel.bind()
    ~~~~~~~~~~~^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/dispatcher.py", line 325, in bind
    cufunc = self._codelibrary.get_cufunc()
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/codegen.py", line 318, in get_cufunc
    cubin = self.get_cubin(cc=device.compute_capability)
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/codegen.py", line 296, in get_cubin
    self._link_all(linker, cc, ignore_nonlto=False)
    ~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/codegen.py", line 263, in _link_all
    linker.add_file_guess_ext(path, ignore_nonlto)
    ~~~~~~~~~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/driver.py", line 2900, in add_file_guess_ext
    self.add_cu(path_or_code.data, path_or_code.name)
    ~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/driver.py", line 2830, in add_cu
    self.add_ptx(ptx.encode(), ptx_name)
    ~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/driver.py", line 3231, in add_ptx
    return self._add_data(enums.CU_JIT_INPUT_PTX, ptx, name)
           ~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/wangm/miniforge3/envs/numbast/lib/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/driver.py", line 3258, in _add_data
    raise LinkerError("%s\n%s" % (e, self.error_log))
numba.cuda.cudadrv.driver.LinkerError: [300] Call to cuLinkAddData results in CUDA_ERROR_INVALID_SOURCE
error   : Prototype doesn't match for 'add' in '<unnamed-cu>.ptx', first defined in '<unnamed-cu>.ptx'

Steps/Code to reproduce bug

import numpy as np
from numba import cuda
from numba.types import float16

src = """
#include <cuda_fp16.h>
extern "C"
__device__ int add(half * res, half a, half b) {
    *res = a + b;
    return 0;
}
"""
add = cuda.declare_device("add", float16(float16, float16), link=cuda.CUSource(src))
@cuda.jit
def fails(arr):
    arr[0] = add(float16(3.14), float16(3.14))


arr = np.zeros(1, dtype=np.float16)
fails[1, 1](arr)
print(arr)

Execute with

NUMBA_CUDA_USE_NVIDIA_BINDING=0 python minrepro.py

Expected behavior
Prints [6.28]

Environment details (please complete the following information):
Conda installed ctk, numba-cuda

[isVoid]

In non-LTO mode, I got the following PTX:

---------------------------------ASSEMBLY fails---------------------------------
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-35583870
// Cuda compilation tools, release 12.8, V12.8.93
// Based on NVVM 7.0.1
//

.version 8.7
.target sm_86
.address_size 64

        // .globl       _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE
.extern .func  (.param .b32 func_retval0) add
(
        .param .b64 add_param_0,
        .param .b32 add_param_1,
        .param .b32 add_param_2
)
;
.common .global .align 8 .u64 _ZN08NumbaEnv8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE;

.visible .entry _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE(
        .param .u64 _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_0,
        .param .u64 _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_1,
        .param .u64 _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_2,
        .param .u64 _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_3,
        .param .u64 _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_4,
        .param .u64 _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_5,
        .param .u64 _ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_6
)
{
        .local .align 2 .b8     __local_depot0[2];
        .reg .b64       %SP;
        .reg .b64       %SPL;
        .reg .pred      %p<3>;
        .reg .b16       %rs<5>;
        .reg .b32       %r<4>;
        .reg .f64       %fd<3>;
        .reg .b64       %rd<5>;


        mov.u64         %SPL, __local_depot0;
        cvta.local.u64  %SP, %SPL;
        ld.param.u64    %rd2, [_ZN8__main__5failsB2v1B92cw51cXTLSUwv1sDUaKthoaNgpqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedE_param_4];
        add.u64         %rd3, %SP, 0;
        add.u64         %rd1, %SPL, 0;
        mov.f64         %fd2, 0d40091EB851EB851F;
        // begin inline asm
        cvt.rn.f16.f64 %rs1, %fd2;
        // end inline asm
        // begin inline asm
        cvt.rn.f16.f64 %rs2, %fd2;
        // end inline asm
        mov.u16         %rs3, 0;
        st.local.u16    [%rd1], %rs3;
        cvt.u32.u16     %r2, %rs1;
        cvt.u32.u16     %r3, %rs2;
        { // callseq 0, 0
        .reg .b32 temp_param_reg;
        .param .b64 param0;
        st.param.b64    [param0+0], %rd3;
        .param .b32 param1;
        st.param.b32    [param1+0], %r2;
        .param .b32 param2;
        st.param.b32    [param2+0], %r3;
        .param .b32 retval0;
        call.uni (retval0), 
        add, 
        (
        param0, 
        param1, 
        param2
        );
        ld.param.b32    %r1, [retval0+0];
        } // callseq 0
        setp.eq.s32     %p1, %r1, 0;
        @%p1 bra        $L__BB0_2;

        setp.ne.s32     %p2, %r1, -2;
        @%p2 bra        $L__BB0_3;

$L__BB0_2:
        cvta.to.global.u64      %rd4, %rd2;
        ld.local.u16    %rs4, [%rd1];
        st.global.u16   [%rd4], %rs4;

$L__BB0_3:
        ret;

}
================================================================================
-----------------------------ASSEMBLY <unnamed-cu>------------------------------
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-35583870
// Cuda compilation tools, release 12.8, V12.8.93
// Based on NVVM 7.0.1
//

.version 8.7
.target sm_86
.address_size 64

        // .globl       add

.visible .func  (.param .b32 func_retval0) add(
        .param .b64 add_param_0,
        .param .align 2 .b8 add_param_1[2],
        .param .align 2 .b8 add_param_2[2]
)
{
        .reg .b16       %rs<4>;
        .reg .b32       %r<2>;
        .reg .b64       %rd<2>;


        ld.param.u64    %rd1, [add_param_0];
        ld.param.u16    %rs2, [add_param_1];
        ld.param.u16    %rs3, [add_param_2];
        // begin inline asm
        {add.f16 %rs1,%rs2,%rs3;
}
        // end inline asm
        st.u16  [%rd1], %rs1;
        mov.u32         %r1, 0;
        st.param.b32    [func_retval0+0], %r1;
        ret;

}

The signature declared for add is

.extern .func  (.param .b32 func_retval0) add
(
        .param .b64 add_param_0,
        .param .b32 add_param_1,
        .param .b32 add_param_2
)
;

But actual add is defined as

.visible .func  (.param .b32 func_retval0) add(
        .param .b64 add_param_0,
        .param .align 2 .b8 add_param_1[2],
        .param .align 2 .b8 add_param_2[2]
)