Fp16 constant weights by mloubout · Pull Request #3015 · devitocodes/devito

mloubout · 2026-09-02T21:09:13Z

No description provided.

codecov · 2026-09-02T21:20:58Z

Codecov Report

❌ Patch coverage is 77.77778% with 6 lines in your changes missing coverage. Please review.
✅ Project coverage is 83.77%. Comparing base (85d54c7) to head (b90c696).

Files with missing lines	Patch %	Lines
devito/operator/operator.py	0.00%	2 Missing and 1 partial ⚠️
devito/ir/cgen/printer.py	80.00%	1 Missing and 1 partial ⚠️
devito/ir/iet/visitors.py	85.71%	0 Missing and 1 partial ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##             main    #3015      +/-   ##
==========================================
- Coverage   83.77%   83.77%   -0.01%     
==========================================
  Files         257      257              
  Lines       54992    55007      +15     
  Branches     4705     4708       +3     
==========================================
+ Hits        46072    46081       +9     
- Misses       8108     8111       +3     
- Partials      812      815       +3

Flag	Coverage Δ
pytest-gpu-aomp-amdgpuX	`68.57% <62.96%> (-0.02%)`	⬇️
pytest-gpu-gcc-	`78.42% <77.77%> (-0.01%)`	⬇️
pytest-gpu-icx-	`78.32% <77.77%> (-0.05%)`	⬇️
pytest-gpu-nvc-nvidiaX	`69.10% <62.96%> (+0.01%)`	⬆️

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

FabioLuporini

IIRC (a may be wrong and confusing it with the other PRO PR) there may have been leftover tiny comments in the old OSS PR that we might be able to address here

FabioLuporini · 2026-09-04T08:19:18Z


    INTERP_MODE = 'direct'
+
+    HALF_ARITH = False


you've got to move it down below -- currently it's right in between "INTERP_MODE" and its docstring

FabioLuporini · 2026-09-04T08:20:15Z


    def _hashable_content(self):
-        return (self.name, self.dimension, str(self.weights), self.scope)
+        # NOTE: `dtype` belongs here. The same coefficients at two precisions


this NOTE can go, it's obvious

FabioLuporini · 2026-09-04T08:26:52Z

+        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)


this seems weird to me:

get_printer(self.printer

FabioLuporini · 2026-09-04T08:27:39Z

+        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)


why, instead of this new private method, don't you just pass self.ccode(obj) and let the printer handle the whole printing of the Array? ie u let it return c.Initializer(...)

FabioLuporini · 2026-09-04T08:28:13Z

    def _print_ListInitializer(self, expr):
        return f"{{{', '.join(self._print(i) for i in expr.params)}}}"

+    def initvalue(self, init, dtype):


if this really is necessary (see other comments below) and if it really needs to be a public method, it should be put among the public methods, not among the private ones

EdCaunt

Mainly comment rephrasing, but the content looks good to me

EdCaunt · 2026-09-07T09:47:54Z


+    HALF_ARITH = False
+    """
+    Whether an Operator working in half precision carries the arithmetic there


Nitpick: "carries out arithmetic"

EdCaunt · 2026-09-07T09:48:54Z

+    HALF_ARITH = False
+    """
+    Whether an Operator working in half precision carries the arithmetic there
+    too, rounding its literals and its FD weights to half. Off by default: half


Comment after "half is a storage format" seems a bit out of context? A Claude-ism I guess?

EdCaunt · 2026-09-07T09:57:37Z

            if not obj._mem_constant or init.is_numeric:
-                value = c.Initializer(value, self.ccode(init))
+                # NOTE: printed at the Array's own precision, not the
+                # Operator's: the two differ for a narrow Array, and it is the


What makes an array "narrow"? From what I can tell it means it contains a single dtype. Maybe just state this plainly?

EdCaunt · 2026-09-07T09:59:10Z

    @cached_property
    def _printer(self):
+        # A Target may offer a second printer for Operators that have opted
+        # into carrying their precision into the arithmetic


Comment could be clearer: "A Target may offer a second printer for Operators using half-precision arithmetic"

The Weights of a non-expanded derivative were always built at the default precision, so a `float16` stencil got `float` coefficients. Every wavefield*weight product then bound to the mixed-precision operators and was promoted, defeating the point of the half-precision wavefield.

`_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.

`_prec` floors an untyped real literal at `float32` so that an integer default doesn't degrade the arithmetic around it. That floor also caught `float16`, which is never a fallback but an explicit request, so every literal in a half-precision Operator printed one type too wide. Only apply the floor when the default is not already a real type.

The stability check sums the whole field and asks whether the result is finite. The accumulator took the field's own dtype, so in half precision it overflowed within a few thousand points and reported an instability that wasn't there -- making `errctl=max`, the very option one reaches for to diagnose a suspected instability, unusable exactly where it is needed. Give it at least single precision.

A real literal in an otherwise integer expression is emitted at the Operator's precision, floored at `float32` so that an integer default does not degrade it. An Operator working in half wants that floor most of the time -- half is a storage format, and the accuracy of the literals is worth more than the width of the multiply -- but not always. Give the printer a flag for it, off by default, and have `_printer` pick up a Target's second printer where one is offered.

The same coefficients at two precisions are two different arrays, but neither `__eq__` nor `_hashable_content` looked at the dtype, so the first one built answered for both. An Operator asking for its weights in one precision would be handed whichever an earlier Operator had cached. Compare and hash on it. The name goes in rather than the type itself, which does not order and so cannot be sorted alongside the rest.

Whether an Operator working in half also computes in half decides what is calculated, not how quickly: the literals and the FD coefficients are rounded to three decimal digits. That is a mathematical choice, so it belongs with `interp-mode` in `sym_opt` rather than among the codegen options, and is validated and defaulted alongside it.

An Array initializer already reached the printer, via `ccode` on a `ListInitializer`; what it did not do was reach it at the Array's own precision. The elements were printed with the Operator's settings, so a narrow Array sitting in an Operator whose arithmetic is left at the default width had its entries emitted at the wider type. Pass the Array's dtype to `ccode`, as `Expression` already does, rather than route the initializer around the printer through an `initvalue` hook of its own. A target that needs to spell its literals differently overrides `_print_ListInitializer`, which is the ordinary extension point. With the precision now correct at the point of printing, `_prec` no longer needs to be told whether the arithmetic was narrowed on purpose: a real literal takes the precision it is being printed at, and the `float32` floor applies only where that is not itself a float. That is the same value as before for every dtype other than `float16`.

Printing an initializer at the Array's dtype was right, but dropping the `float32` floor from `_prec` to get there was not: `Expression` also passes a dtype, and there it is the width the expression operates at, which does not constrain the width of a literal within it. A half Operator's updates were narrowing their literals with no opt-in. Restore the floor and let a caller waive it with `exact_prec`, which the Array initializer sets: its element type is fixed, so a literal that does not fit is ill-typed rather than merely less accurate.

mloubout added the API api (symbolics, types, ...) label Sep 2, 2026

mloubout force-pushed the fp16-constant-weights branch from 75d6b31 to 9759188 Compare September 3, 2026 12:10

FabioLuporini reviewed Sep 4, 2026

View reviewed changes

mloubout force-pushed the fp16-constant-weights branch 2 times, most recently from 9f3c81d to 38dbac3 Compare September 4, 2026 14:17

mloubout added the no-pro-trigger Skip the devitopro submodule update on merge label Sep 4, 2026

FabioLuporini approved these changes Sep 7, 2026

View reviewed changes

EdCaunt reviewed Sep 7, 2026

View reviewed changes

mloubout force-pushed the fp16-constant-weights branch 2 times, most recently from 76d665c to 1793f4b Compare September 8, 2026 00:27

mloubout added 8 commits September 8, 2026 06:29

mloubout force-pushed the fp16-constant-weights branch from 1793f4b to 7597bee Compare September 8, 2026 10:30

mloubout merged commit 33c3b12 into main Sep 8, 2026
42 checks passed

mloubout deleted the fp16-constant-weights branch September 8, 2026 23:49

Navigation Menu

Uh oh!

Fp16 constant weights - #3015

Fp16 constant weights#3015
mloubout merged 9 commits into
mainfrom
fp16-constant-weights

mloubout commented Sep 2, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Sep 2, 2026 •

edited

Loading

Uh oh!

FabioLuporini left a comment

Uh oh!

FabioLuporini Sep 4, 2026

Uh oh!

FabioLuporini Sep 4, 2026

Uh oh!

FabioLuporini Sep 4, 2026

Uh oh!

FabioLuporini Sep 4, 2026

Uh oh!

FabioLuporini Sep 4, 2026

Uh oh!

EdCaunt left a comment

Uh oh!

EdCaunt Sep 7, 2026

Uh oh!

EdCaunt Sep 7, 2026

Uh oh!

EdCaunt Sep 7, 2026

Uh oh!

EdCaunt Sep 7, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

mloubout commented Sep 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov Bot commented Sep 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

FabioLuporini left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

EdCaunt left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

mloubout commented Sep 2, 2026 •

edited

Loading

codecov Bot commented Sep 2, 2026 •

edited

Loading