text: Title: [numba.cuda] StructModel + FP16 fails with CUDA because make_attribute_wrapper assumes default_manager 
  Body: This
  ```python
  from numba import cuda
  from numba import types
  from numba.extending import models, make_attribute_wrapper, lower_builtin, as_numba_type, type_callable
  from numba.cuda.models import register_model
  from numba.core import cgutils
  import numba.cuda.cudaimpl
  
  class float16x2(object):
      pass
  
  class float16x2Type(types.Type):
      def __init__(self):
          super(float16x2Type, self).__init__(name='float16x2')
  
  float16x2_type = float16x2Type()
  
  ## Type inference (Python --> FE type)
  
  as_numba_type.register(float16x2, float16x2_type)
  
  @type_callable(float16x2)
  def type_float16x2(context):
      def typer(x, y):
          if all(isinstance(v, types.Float) for v in [x, y]):
              return float16x2_type
      return typer
  
  ### Lowering (FE types --> LLVM)
  
  # Data model
  @register_model(float16x2Type)
  class float16x2TypeModel(models.StructModel):
      def __init__(self, dmm, fe_type):
          members = [(f, types.float16) for f in ['x', 'y']]
          models.StructModel.__init__(self, dmm, fe_type, members)
  
  for field in ['x', 'y']:
      make_attribute_wrapper(float16x2Type, field, field)
  
  # Constructors
  @lower_builtin(float16x2, types.Float, types.Float)
  def impl(context, builder, sig, args):
      typ = sig.return_type
      obj = cgutils.create_struct_proxy(typ)(context, builder)
      x, y = args
      x = numba.cuda.cudaimpl.float_to_float16_cast(context, builder, sig.args[0], types.float16, x)
      y = numba.cuda.cudaimpl.float_to_float16_cast(context, builder, sig.args[1], types.float16, y)
      obj.x, obj.y = x, y
      return obj._getvalue()
  
  @cuda.jit
  def f():
      a = float16x2(1.0, 2.0)
      print(a.y)
  
  f[1, 1, 0, 0]()
  cuda.synchronize()
  ```
  fails with
  ```
  numba.core.errors.TypingError: Failed in cuda mode pipeline (step: nopython frontend)
  Internal error at resolving type of attribute "y" of "a".
  <class '__main__.float16x2Type'>
  During: typing of get attribute at [redacted] (54)
  Enable logging at debug level for details.
  
  File "[redacted]", line 54:
  def f():
      <source elided>
      a = float16x2(1.0, 2.0)
      print(a.y)
      ^
  ```
  
  The reason seems to be that `make_attribute_wrapper` [here](https://github.com/numba/numba/blob/main/numba/core/extending.py#L279C38-L279C53) uses
  ```
  from numba.core.datamodel import default_manager
  ```
  and, as such, assumes the default data model. This means for FP16 it will uses this generic data model [here](https://github.com/numba/numba/blob/7b7137a1f2bf003e8aa42255f571d0947654e7ce/numba/core/datamodel/models.py#L363) instead of the CUDA specific one from [here](https://github.com/numba/numba/blob/7b7137a1f2bf003e8aa42255f571d0947654e7ce/numba/cuda/models.py#L35)
  
  cc @gmarkall 
  