compiler: Print an Array initializer at the Array's own precision

`_gen_value` printed the initializer with the printer's default dtype
rather than the Array's, which stamped a `float` suffix onto the entries
of a `double` Array and silently rounded them to single precision.

Route it through a new `initvalue` printer hook, which also gives the
targets a place to specialize an initializer whose type cannot be built
from a plain literal.

Author: mloubout

devito/ir/cgen/printer.py

@@ -373,6 +373,17 @@ def _print_FieldFromComposite(self, expr):
     def _print_ListInitializer(self, expr):
         return f"{{{', '.join(self._print(i) for i in expr.params)}}}"
 
+    def initvalue(self, init, dtype):
+        """
+        Print the aggregate initializer `init` of an Array of type `dtype`.
+
+        Kept separate from `_print_ListInitializer` because a static
+        initializer, unlike an expression, cannot rely on implicit conversions:
+        some types (e.g. CUDA's `__half`) are only constructible from a literal
+        via a runtime call, which is illegal in that position.
+        """
+        return self._print(init)
+
     def _print_IndexedPointer(self, expr):
         base = self._print(expr.base)
         return f"{base}{''.join(f'[{self._print(i)}]' for i in expr.index)}"

devito/ir/iet/visitors.py

@@ -360,12 +360,21 @@ def _gen_value(self, obj, mode=1, masked=()):
         if obj.is_Array and obj.initvalue is not None and mode == 1:
             init = ListInitializer(obj.initvalue)
             if not obj._mem_constant or init.is_numeric:
-                value = c.Initializer(value, self.ccode(init))
+                value = c.Initializer(value, self._gen_initvalue(obj, init))
         elif obj.is_LocalObject and obj.initvalue is not None and mode == 1:
             value = c.Initializer(value, self.ccode(obj.initvalue))
 
         return value
 
+    def _gen_initvalue(self, obj, init):
+        """
+        Convert the aggregate initializer `init` of the Array `obj` into a C
+        string, delegating to the printer so that languages whose types cannot
+        be built from plain literals (e.g. CUDA's `__half`) can specialize it.
+        """
+        printer = get_printer(self.printer, obj.dtype)
+        return printer.initvalue(init, obj.dtype)
+
     def _gen_rettype(self, obj):
         try:
             return self._gen_value(obj, 0).typename