Add dace::float32sr type to DaCe #2148

creavin · 2025-09-14T20:33:33Z

This new type enables DaCe users to perform calculations with stochastic rounding in single precision.

This change is validated with unit tests.

phschaad · 2025-10-30T09:19:40Z

It appears this PR is still failing tests, will there be any more movement on this?

acalotoiu

The tests would need to be fixed. I am also not sure if we should aim merging to main before trying to work on improving the performance of the SR implementation.

creavin · 2025-11-17T11:52:33Z

Ack, will (attempt to) fix the tests this week

This new type enables DaCe users to perform calculations with stochastic rounding in single precision. This change is validated with unit tests.

creavin · 2025-11-23T12:51:03Z

Okay, all tests pass except for the FPGA Tests. I fixed the linking error. Now it fails because I use an incompatible abs function and it doesn't support thread-local [1].

I don't have a setup in which I can debug this. I see Lex's comment that the performance of this should be evaluated further before pulling it into main. What I would suggest is that maybe you move this to a feature / experimental branch and when this eval is being done, one tries run this on your FPGA setup and find the compatible alternatives to abs and thread-local.

[](https://github.com/spcl/dace/actions/runs/19598553709/job/56126669713?pr=2148#step:4:446)WARNING: [HLS 207-5324] using integer absolute value function 'abs' when argument is of floating point type (/home/timos/actions-runner/_work/dace/dace/dace/codegen/../runtime/include/dace/stocastic_rounding.h:58:13)

[](https://github.com/spcl/dace/actions/runs/19598553709/job/56126669713?pr=2148#step:4:447)INFO: [HLS 207-4500] use function 'std::abs' instead (/home/timos/actions-runner/_work/dace/dace/dace/codegen/../runtime/include/dace/stocastic_rounding.h:58:13)

[](https://github.com/spcl/dace/actions/runs/19598553709/job/56126669713?pr=2148#step:4:448)ERROR: [HLS 207-3638] thread-local storage is not supported for the current target (/home/timos/actions-runner/_work/dace/dace/dace/codegen/../runtime/include/dace/stocastic_rounding.h:196:7)

tbennun · 2026-01-10T05:05:51Z

@creavin Maybe try again after #2252 is merged.

creavin · 2026-01-10T14:50:50Z

Ack, will do

…

On Sat 10 Jan 2026, 06:06 Tal Ben-Nun, ***@***.***> wrote: *tbennun* left a comment (spcl/dace#2148) <#2148 (comment)> @creavin <https://github.com/creavin> Maybe try again after #2252 <#2252> is merged. — Reply to this email directly, view it on GitHub <#2148 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/ABD4ULXZ63Y6KMIZI5IOEHT4GCCEJAVCNFSM6AAAAACGPQOZBKVHI2DSMVQWIX3LMV43OSLTON2WKQ3PNVWWK3TUHMZTOMZRHAZTOMJSGY> . You are receiving this because you were mentioned.Message ID: ***@***.***>

phschaad

Some comments and minor changes before I can accept, but overall happy with the change, thanks!

phschaad · 2026-01-13T09:41:11Z

dace/frontend/python/replacements/operators.py

        restype = np_result_type(dtypes_for_result)
        coarse_result_type = None
-        if result_type in complex_types:
+        if restype in complex_types:


What is/was result_type, why did this change? Was this a bug?

This is a bug introduced by this refactor: c909f8b#diff-b20441227a628465d1c6ca8819915b77cf036f8fc1a8e26fbceb8a930fde9d1dR353

The origional code, formerly in the replacements.py file, looked like

else: # Operators with 3 or more arguments result_type = _np_result_type(dtypes_for_result) coarse_result_type = None if result_type in complex_types: coarse_result_type = 3 # complex elif result_type in float_types: coarse_result_type = 2 # float elif result_type in signed_types: coarse_result_type = 1 # signed integer, bool else: coarse_result_type = 0 # unsigned integer for i, t in enumerate(coarse_types): if t != coarse_result_type: casting[i] = _cast_str(result_type)

I believe this var was renamed to restype to not conflict when the "result_type" function

phschaad · 2026-01-13T09:41:52Z

dace/libraries/blas/nodes/dot.py

-            raise TypeError(f"Data types of input and output must be equal: {desc_x.dtype}, {desc_res.dtype}")
+            input_types = (desc_x.dtype.base_type, desc_res.dtype.base_type)
+            if dace.float32sr in input_types and dace.float32sr in input_types:
+                pass  # ignore mismatch if it is stochastically rounded


I'm not an expert on SR types, so my question here is: can we safely do that, or rather - why can we safely do that?

I have made a typo here -- thank you for catching that. Line 210 should be:
if dace.float32sr in input_types and dace.float32 in input_types:. I will also rename input_types to arg_types and add more comments

What this does and why it's needed: when using stocastic rounding, a legitimate (i.e not a bug) mismatch between the input and output arguments may arise where one argument is a float32sr and the other is a float32 (round-to-nearest). The underlying data type is the same so this should not cause the validation to fail. This is a check for that edge-case

dace/runtime/include/dace/stocastic_rounding.h

phschaad · 2026-01-13T09:44:32Z

dace/runtime/include/dace/stocastic_rounding.h

+    #endif
+
+    DACE_HOST_DEVICE static float stochastic_round(double x) {
+        // uint64_t rbits = lcg64();


Do these commented out lines still carry significance?

So as part of my thesis, I make a comparison between five different rounding modes. I kept the functions and these commented out invocations partly for posterity and partly to allow future devs to play around with it.

I think there is value in keeping this. I will make this into one compact comment to reduce its footprint. If you disgree, I can remove the comment and the associated functions

tests/codegen/float32sr.py

creavin · 2026-01-13T09:48:56Z

Super, I'm happy to make them. Will get back to you soon

…

On Tue 13 Jan 2026, 10:46 Philipp Schaad, ***@***.***> wrote: ***@***.**** requested changes on this pull request. Some comments and minor changes before I can accept, but overall happy with the change, thanks! ------------------------------ In dace/frontend/python/replacements/operators.py <#2148 (comment)>: > @@ -336,18 +370,20 @@ def result_type(arguments: Sequence[Union[str, Number, symbolic.symbol, sp.Basic else: # Operators with 3 or more arguments restype = np_result_type(dtypes_for_result) coarse_result_type = None - if result_type in complex_types: + if restype in complex_types: What is/was result_type, why did this change? Was this a bug? ------------------------------ In dace/libraries/blas/nodes/dot.py <#2148 (comment)>: > @@ -204,7 +206,11 @@ def validate(self, sdfg, state): if desc_x.dtype != desc_y.dtype: raise TypeError(f"Data types of input operands must be equal: {desc_x.dtype}, {desc_y.dtype}") if desc_x.dtype.base_type != desc_res.dtype.base_type: - raise TypeError(f"Data types of input and output must be equal: {desc_x.dtype}, {desc_res.dtype}") + input_types = (desc_x.dtype.base_type, desc_res.dtype.base_type) + if dace.float32sr in input_types and dace.float32sr in input_types: + pass # ignore mismatch if it is stochastically rounded I'm not an expert on SR types, so my question here is: can we safely do that, or rather - why can we safely do that? ------------------------------ In dace/runtime/include/dace/stocastic_rounding.h <#2148 (comment)>: > @@ -0,0 +1,390 @@ +#ifndef __DACE_SROUND_H Please include the appropriate copyright comment ------------------------------ In dace/runtime/include/dace/stocastic_rounding.h <#2148 (comment)>: > +#include <cstdint> +#include <cstring> +#include <iomanip> +#include <iostream> +#include <limits> +#include <random> +#include <thread> + +#ifdef __CUDACC__ +#define DACE_HOST_DEVICE __host__ __device__ +#include <curand_kernel.h> +#else +#define DACE_HOST_DEVICE +#endif + +namespace dace { This file could in general use some more documenting comments ------------------------------ In dace/runtime/include/dace/stocastic_rounding.h <#2148 (comment)>: > + uint64_t& rng_state = get_rng_state_64(); + rng_state = rng_state * 6364136223846793005ULL + 1442695040888963407ULL; // 64-bit LCG constants + return rng_state; + } + + DACE_HOST_DEVICE static inline uint64_t xorshift64() { + uint64_t& rng_state = get_rng_state_64(); + rng_state ^= rng_state << 13; + rng_state ^= rng_state >> 7; + rng_state ^= rng_state << 17; + return rng_state; + } + #endif + + DACE_HOST_DEVICE static float stochastic_round(double x) { + // uint64_t rbits = lcg64(); Do these commented out lines still carry significance? ------------------------------ In tests/codegen/float32sr.py <#2148 (comment)>: > + upper = rounded_val + + print(f"lower: {lower}, upper: {upper}") + + print(f" C: {higher_prec_val}") + lower_diff = abs(higher_prec_val - lower) + print(f"LF: {lower_diff}") + upper_diff = abs(upper - higher_prec_val) + print(f"UF: {upper_diff}") + theoretical_prob = lower_diff / (lower_diff + upper_diff) + + return lower_diff, upper_diff, theoretical_prob + + +if __name__ == "__main__": + import pytest Please check the other test files on how we run tests from the main function. Note that this is done for debugging purposes and thus is important. — Reply to this email directly, view it on GitHub <#2148 (review)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/ABD4ULS7CE4PPNBJ6OBZ5G34GS5GPAVCNFSM6AAAAACGPQOZBKVHI2DSMVQWIX3LMV43YUDVNRWFEZLROVSXG5CSMV3GSZLXHMZTMNJUHA2TSMRRHE> . You are receiving this because you were mentioned.Message ID: ***@***.***>

creavin

Addressed all feedback item. Corrections to follow in next commit

creavin · 2026-01-18T16:24:56Z

dace/frontend/python/replacements/operators.py

        restype = np_result_type(dtypes_for_result)
        coarse_result_type = None
-        if result_type in complex_types:
+        if restype in complex_types:


This is a bug introduced by this refactor: c909f8b#diff-b20441227a628465d1c6ca8819915b77cf036f8fc1a8e26fbceb8a930fde9d1dR353

The origional code, formerly in the replacements.py file, looked like

else: # Operators with 3 or more arguments result_type = _np_result_type(dtypes_for_result) coarse_result_type = None if result_type in complex_types: coarse_result_type = 3 # complex elif result_type in float_types: coarse_result_type = 2 # float elif result_type in signed_types: coarse_result_type = 1 # signed integer, bool else: coarse_result_type = 0 # unsigned integer for i, t in enumerate(coarse_types): if t != coarse_result_type: casting[i] = _cast_str(result_type)

I believe this var was renamed to restype to not conflict when the "result_type" function

dace/runtime/include/dace/stocastic_rounding.h

creavin · 2026-01-18T16:56:11Z

dace/libraries/blas/nodes/dot.py

-            raise TypeError(f"Data types of input and output must be equal: {desc_x.dtype}, {desc_res.dtype}")
+            input_types = (desc_x.dtype.base_type, desc_res.dtype.base_type)
+            if dace.float32sr in input_types and dace.float32sr in input_types:
+                pass  # ignore mismatch if it is stochastically rounded


I have made a typo here -- thank you for catching that. Line 210 should be:
if dace.float32sr in input_types and dace.float32 in input_types:. I will also rename input_types to arg_types and add more comments

What this does and why it's needed: when using stocastic rounding, a legitimate (i.e not a bug) mismatch between the input and output arguments may arise where one argument is a float32sr and the other is a float32 (round-to-nearest). The underlying data type is the same so this should not cause the validation to fail. This is a check for that edge-case

dace/runtime/include/dace/stocastic_rounding.h

creavin · 2026-01-18T17:15:29Z

dace/runtime/include/dace/stocastic_rounding.h

+    #endif
+
+    DACE_HOST_DEVICE static float stochastic_round(double x) {
+        // uint64_t rbits = lcg64();


So as part of my thesis, I make a comparison between five different rounding modes. I kept the functions and these commented out invocations partly for posterity and partly to allow future devs to play around with it.

I think there is value in keeping this. I will make this into one compact comment to reduce its footprint. If you disgree, I can remove the comment and the associated functions

tests/codegen/float32sr.py

phschaad

LGTM, thanks for addressing the comments

creavin assigned acalotoiu Sep 14, 2025

creavin requested a review from acalotoiu September 14, 2025 20:34

creavin unassigned acalotoiu Sep 14, 2025

phschaad added the waiting for author Will be closed in 30 days label Oct 30, 2025

acalotoiu reviewed Nov 17, 2025

View reviewed changes

creavin added 3 commits November 22, 2025 14:35

Add dace::float32sr to DaCe

c2de579

This new type enables DaCe users to perform calculations with stochastic rounding in single precision. This change is validated with unit tests.

Apply linting

fd83b4b

Fix gemm cast and dace result cast

d0ee5c6

creavin force-pushed the single-prec-sr branch from 1c51e8a to d0ee5c6 Compare November 22, 2025 13:59

creavin added 2 commits November 22, 2025 16:06

Fix gemm cast arg placement

3c33f58

Try inline the rng for fpga compat

ac7aa7a

creavin removed the waiting for author Will be closed in 30 days label Nov 23, 2025

creavin requested a review from phschaad November 23, 2025 12:52

Merge branch 'main' into single-prec-sr

31ac0c8

phschaad requested changes Jan 13, 2026

View reviewed changes

creavin commented Jan 20, 2026

View reviewed changes

creavin added 2 commits January 20, 2026 12:06

Minor bug fix; Commented SR code

6fd0397

Merge branch 'main' into single-prec-sr

6178ed9

creavin requested a review from phschaad January 20, 2026 12:40

phschaad approved these changes Jan 20, 2026

View reviewed changes

phschaad added this pull request to the merge queue Jan 20, 2026

Merged via the queue into main with commit c9576df Jan 20, 2026
12 checks passed

phschaad deleted the single-prec-sr branch January 20, 2026 15:22

Add dace::float32sr type to DaCe #2148

Add dace::float32sr type to DaCe #2148

Uh oh!

Conversation

creavin commented Sep 14, 2025

Uh oh!

phschaad commented Oct 30, 2025

Uh oh!

acalotoiu left a comment

Choose a reason for hiding this comment

Uh oh!

creavin commented Nov 17, 2025

Uh oh!

creavin commented Nov 23, 2025

Uh oh!

tbennun commented Jan 10, 2026

Uh oh!

creavin commented Jan 10, 2026 via email

Uh oh!

phschaad 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!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

creavin commented Jan 13, 2026 via email

Uh oh!

creavin left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

phschaad left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants