[MRG] Float32 support #981

mstimberg · 2018-07-12T08:36:48Z

Due to summer vacations this will still take a while to merge I think, but I'm putting it out here since it is basically ready from my side and it should be of high interest to the @brian-team/brian2cuda and @brian-team/brian2genn backends.

There are a few tests (mostly the multicompartmental stuff) that test numerical results vs. theoretical predictions that are just off too far with single precision, so I decided to simply skip them when single precision is used. In the long run, we should also probably split these type of tests from the more basic unit/integration tests.

Travis and appveyor run tests with single precision (but only for Python 2.7, the test suits already take way too long...) and everything seems to work fine.

A remark about the clock: The variables t and dt always use double precision regardless of the preference setting, because everything gets very imprecise otherwise. To avoid problems, other variables or subexpressions that are direct functions of time (e.g. lastspike) are using the same dtype.

Fixes #417

…float64)

… dtype

…loating point variables and lower required accuracy in some cases

…ult float type

# Conflicts: # brian2/tests/test_functions.py # brian2/tests/test_refractory.py # brian2/tests/test_subgroup.py # brian2/tests/test_timedarray.py

the same dtype as time)

denisalevi · 2018-08-09T16:29:43Z

Hi @mstimberg. Does brian2genn already work out of the box with this brian2 preference by any chance? :) or would it be possible to add it for our comparisons (I imagine it to be easy to translate the brian pref to genn pref?)

mstimberg · 2018-08-13T16:30:16Z

No, it doesn't work with Brian2GeNN (it's not even merged in Brian 2 yet :) ). In principle, it should be easy to translate, but I see potential problems with variables like lastupdate that we force to be double precision in Brian.

thesamovar

This looks good to me except for the doctest issue which would be good to address if we can?

I'm actually amazed by how few changes were needed to the actual code (not to the tests) to make this work. Great work!

thesamovar · 2018-08-24T12:05:59Z

brian2/core/core_preferences.py

@@ -14,16 +14,14 @@ def dtype_repr(dtype):


 def default_float_dtype_validator(dtype):
-    return dtype is float64
+    return dtype in [float32, float64]


I don't think we need it as a priority, but I wonder how much work it would be to allow for other types such as float16 or float128?

I don't think we'd need any additional changes for that, it's just more effort for testing. But we could certainly think about it again when a use case arises.

thesamovar · 2018-08-24T12:18:02Z

brian2/monitors/spikemonitor.py

@@ -237,16 +237,18 @@ def values(self, var):
        Examples
        --------
        >>> from brian2 import *
-        >>> G = NeuronGroup(2, """dv/dt = 100*Hz : 1
+        >>> G = NeuronGroup(2, """dv/dt = 100.0001*Hz : 1


This is really ugly because it's actually in the documentation. Is there no way around this?

I'll have a look, we can certainly come up with a different kind of example.

There are quite a few examples like this (not just this one).

thesamovar · 2018-08-24T12:18:53Z

brian2/monitors/spikemonitor.py

-        array([ 0.5,  0.5,  0.5,  0.5])
-        >>> v_values[1]
-        array([ 1.,  1.])
+        >>> np.set_printoptions(precision=4)  # show fewer digits than default


Same here. I hate to have stuff like this actually in our docs.

thesamovar · 2018-08-24T12:21:44Z

brian2/tests/test_GSL.py

-        assert not all(trace_conventional.v[0]==trace_GSL.v[0]), \
-                ('output of GSL stateupdater is exactly the same as Brians stateupdater (unlikely to be right)')
+        # assert not all(trace_conventional.v[0]==trace_GSL.v[0]), \
+        #         ('output of GSL stateupdater is exactly the same as Brians stateupdater (unlikely to be right)')


Why is this commented out?

Ah, right. The test was a bit weird, checking that the results of two state updaters for the same equations were not exactly the same. It was more a "during-development" test, Charlee used it to make sure that the simulation was actually using her state updater. I think this test failed with float32, i.e. the results were exactly the same. We can simply delete it.

thesamovar · 2018-08-24T12:24:34Z

brian2/tests/test_neurongroup.py

@@ -654,7 +655,7 @@ def test_threshold_reset():
                    threshold='v > 1', reset='v=0.5')
    G.v = np.array([0, 1, 2])
    run(defaultclock.dt)
-    assert_equal(G.v[:], np.array([0, 1, 0.5]))
+    assert_allclose(G.v[:], np.array([0, 1, 0.5]))


Comment not about this particular instance, but I wonder if we lose anything by weakening our tests like this? My feeling is probably not, but it makes me mildly anxious. I guess you've already pretty much covered this with the float-type dependent version of allclose.

In general I think we need to do a better job separating our tests for basic functionality from the tests for mathematical/numerical correctness. Currently, we only have a few tests (e.g. for STDP) that compare the exact results for non-trivial simulations.

Worth creating an issue for? I'm not entirely sure it's necessary, I feel like allclose is probably as good as equal practically. Can we think of an example where it might be a problem that it passes allclose but not equal?

thesamovar · 2018-08-24T12:32:33Z

Oh, by the way do the CI tests now run for float32 and float64?

mstimberg · 2018-08-24T12:51:44Z

Oh, by the way do the CI tests now run for float32 and float64?

They do, or almost: we are testing float32 on all codegen targets and on all platforms, but only for Python 2 (our test suite is already huge as it is...)

thesamovar · 2018-08-24T13:48:51Z

If I just run_nose_tests.py locally will I get float64? How can I run with float32 locally?

mstimberg · 2018-08-24T13:53:35Z

If I just run_nose_tests.py locally will I get float64? How can I run with float32 locally?

By default it uses float64, to get float32 you'll have to pass float_dtype=np.float32 to brian2.test

…nt issues

mstimberg · 2018-08-27T16:28:26Z

Sorry, I have to leave... Can you check that you have the line I linked in my previous comment? Just before running the "codegen-independent" tests it sets set_device('runtime'), therefore I don't understand why it tries to run the tests with standalone.

thesamovar · 2018-08-27T16:31:52Z

That line is there, yep. Might need a bit of help debugging this one. Will be around tomorrow afternoon.

mstimberg · 2018-08-28T14:15:44Z

For debugging this, can you see whether you see the same problem when you run the run_nose_tests_long_and_standalone.py but replace:

success = [brian2.test(long_tests=True, test_standalone='cpp_standalone', float_dtype=np.float32),
           brian2.test(long_tests=True, test_standalone='cpp_standalone', float_dtype=np.float64)]

by

success = [brian2.test([], test_standalone='cpp_standalone', float_dtype=np.float32),
           brian2.test([], test_standalone='cpp_standalone', float_dtype=np.float64)]

If yes, that should at least make the debugging much faster!

tnowotny · 2018-08-28T14:24:47Z

About the GeNN question: We allow variables to be declared as "scalar x" and then x gets the type that was selected with the preference setPrecision(GENN_FLOAT) etc. Any type that is explicitly set to float or double will remain what it is set to. At least this is the intended behaviour. I won't vouch for these mechanisms to have undergone thorough testing.

mstimberg · 2018-08-28T14:45:13Z

Thanks for clearing this up @tnowotny . I don't see us using scalar x anywhere in the Brian2GeNN code, so I think it should work just fine even without us setting the precision. Still, it makes sense to set it to avoid confusion.

thesamovar · 2018-08-28T15:37:28Z

Yep that works! Will open brian dev gitter if you want to debug together.

mstimberg · 2018-08-28T15:52:04Z

Will open brian dev gitter if you want to debug together.

I'm in the gitter chat (not sure whether it's the right one :) )

…ce/reset_device

tnowotny · 2018-08-29T09:42:57Z

However, I just remembered a couple of other uses of the setPrecision information:

we replace literal constants with their appropriate single precision float versions in model code, e.g. 3.1415 becomes 3.1415f
we replace math function names with explicit precision variants, e.g. sin() becomes sinf()
the global variable t that handles time (even though as derived from a long integer variable iT that does the actual counting) is defined according to the current precision.

All of these make it a good idea to set the precision.

thesamovar

Tests now pass!

mstimberg · 2018-08-29T11:01:38Z

Tests now pass!

Same here, I'll merge it right away even though the appveyor tests did not pass yet (this will take > 5h still...)

mstimberg · 2018-08-29T11:41:14Z

the global variable t that handles time (even though as derived from a long integer variable iT that does the actual counting) is defined according to the current precision.

This is a problem, unfortunately, since in Brian t is always stored as a double variable... I don't think we will change this, since t is so ubiqutous in user expressions and with single precision, stuff like mon.v[mon.t>=1*ms] will often not do what the user thinks it should do (since 1*ms is a double). Even just for display, things get ugly:

>>> Quantity(1e-3, dtype=np.float32, dim=second.dim)
1.00000005 * msecond

I see two solutions:

always use GENN_DOUBLE as the precision for GeNN -- individual variables will still have the dtype that Brian defines
fix Brian2GeNN so that it can work with GeNN's single precision t

(1) is trivial to do, but of course we lose most of the performance benefits when every function call etc. still uses double precision.

I'll look into (2), maybe it is not that difficult to fix. We can still raise a warning to make users aware that Brian2GeNN is not calculating exactly the same thing as Brian.

tnowotny · 2018-08-29T11:47:43Z

I can see the logic of using double precision for t. We made the "clean" float version originally for older GPUs which did not support doubles at all. In that case we had to use float for everything. I suspect today that is pure history - opinions @neworderofjamie?
I think solution (1) would be quite counter-productive even though I am also not sure how the different compilers handle calls to sin(x) where x is a single precision float. Do they auto-cast x to double or resolve to a float precision sin function "C++-style"?

neworderofjamie · 2018-08-29T11:53:54Z

Woah this is prescient! In the past (genn-team/genn#183) we talked about abandoning floating point time altogether and I started implementing that. However it totally breaks backwards compatibility and ends up making the model code really ugly so I was thinking that adding an option to make time double precision is probably a pragmatic solution - single precision t gets totally mangled after running my STDP model for 2000s 😢

mstimberg · 2018-08-29T11:59:34Z

I think solution (1) would be quite counter-productive even though I am also not sure how the different compilers handle calls to sin(x) where x is a single precision float. Do they auto-cast x to double or resolve to a float precision sin function "C++-style"?

C++11 (which I think is required by GeNN?) has overloaded math functions for single precision.

neworderofjamie · 2018-08-29T12:01:42Z

Sadly I think the CUDA maths functions are C-style - as in the float ones have an f at the end rather than relying on overloading

tnowotny · 2018-08-29T12:03:27Z

hum ... @neworderofjamie, how were you thinking to make a double precision time version? We could introduce a new GeNN flag for it and then just make t double precision based on it - something like that?

tnowotny · 2018-08-29T12:04:10Z

Brian2GeNN then could just set that GeNN flag and of we would go ... I am probably missing something.

mstimberg · 2018-08-29T12:06:12Z

This sounds good to me. But a general question: is there a way to find out about the GeNN version "at runtime"? Even a simple thing like a version.txt file in the GeNN directory would be enough. This way we could check from Brian2GeNN what GeNN version is used and only set an option if it is supported, warn the user otherwise, etc.

neworderofjamie · 2018-08-29T12:08:36Z

That would be exactly my suggestion

…

On Wed, 29 Aug 2018, 13:04 Thomas Nowotny, ***@***.***> wrote: Brian2GeNN then could just set that GeNN flag and of we would go ... I am probably missing something. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#981 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AGeoGqr6gAM4zpTzIOMMKpadVVJMtlvcks5uVoNFgaJpZM4VMbjM> .

tnowotny · 2018-08-29T12:11:11Z

@neworderofjamie, that is very simple to do. I guess we would rely on the compiler to auto-cast whenever t is used in some simulation code, but may not be the worst thing.
I guess we should do it.
@mstimberg about the version: I don't think we have anything in place but sounds like a really sensible thing to have.

denisalevi · 2018-08-29T12:18:31Z

Sadly I think the CUDA maths functions are C-style - as in the float ones have an f at the end rather than relying on overloading

As far as I know, CUDA math functions are overloaded for single and double precision. At least in CUDA 9.0 (no idea about pervious versions). See here. Even though there is only an example for log given. Correct me if I'm wrong.

tnowotny · 2018-08-29T12:28:39Z

Looks like it in the documentation you referenced. In that case we could relax with our ensureFtype logic and not modify function names for single precision (at least from CUDA 9 onwards).
Independently, this would make @mstimberg's option 1) viable as well (at least from CUDA 9 onwards).

neworderofjamie · 2018-08-29T12:42:02Z

Sadly, for reasons best known to NVIDIA, that seems to only apply to log - see https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__SINGLE.html#group__CUDA__MATH__SINGLE

neworderofjamie · 2018-08-29T12:45:02Z

Actually, apologies, that is just terrible documentation! In maths_functions.hpp there's lots of the following:

__host__ __device__ __cudart_builtin__ float     sin(float in)           { return sinf(in); }

I will check older CUDA's and see if this is consistent - I would IMAGINE it came in when CUDA started supporting C++.

mstimberg · 2018-08-29T12:45:51Z

I don't think that's the case, their documentation is just... not that good. The log function is just an example (on the page you linked, there's also logf, and the double precision docs also pretend that log is double-only).

It seems that this has been in CUDA for a long time, I've found "Introduction to CUDA" slides from 2011 that mention that math functions are overloaded.

mstimberg · 2018-08-29T12:54:39Z

Let's continue discussions about float32 support in Brian2GeNN here: brian-team/brian2genn#69

mstimberg added 19 commits April 13, 2018 17:26

Support using float32 as the default dtype

97ae40c

Fixes #417

Fix a number of doctests to work with float32

4b99db6

Monitors should record t with the dtype it uses in the clock (i.e. …

0d9468f

…float64)

The lastspike/lastupdate variables should use the same dtype as t

5f7d5eb

SpikeGeneratorGroup: period and spike times should be stored with t's…

397d4a6

… dtype

Replace assert_equal with assert_allclose in tests that compare f…

b05380a

…loating point variables and lower required accuracy in some cases

Add an assert_allclose function with tolerance relative to the defa…

f34fb05

…ult float type

Adapt GSL to work with single precision floats

349fb00

Fix a few tests

4d72c40

Fix an issue for dimensionless variables

319104e

Skip tests that need double precision when preference is set to single

66a662f

Merge branch 'master' into float32_support

a4caaf7

# Conflicts: # brian2/tests/test_functions.py # brian2/tests/test_refractory.py # brian2/tests/test_subgroup.py # brian2/tests/test_timedarray.py

Fix a missing import in the tests

676f2d6

Reset numpy's print options in doctests after changing them

8dd5dd8

Add an option to change the default dtype for testing

ea67588

Fix a few doc tests to not use too much precision for their output

3cc424c

Make the refractory test more robust (time-related variables should use

2c1d704

the same dtype as time)

Test 32bit float dtype on the CI servers

2a6e90b

Use assert_allclose with less strict error criterion in a test

c4f2820

moritzaugustin mentioned this pull request Jul 12, 2018

Single precision preference brian-team/brian2cuda#37

Closed

thesamovar reviewed Aug 24, 2018

View reviewed changes

Also test 32bit floating point type on OS X

9431170

mstimberg added 2 commits August 24, 2018 15:55

Change SpikeMonitor/EventMonitor doctests independent of floating poi…

8cda19f

…nt issues

Remove unnecessary test in GSL

885814a

Get the "previous devices" into a known state before testing set_devi…

9732bd1

…ce/reset_device

thesamovar approved these changes Aug 29, 2018

View reviewed changes

mstimberg mentioned this pull request Aug 29, 2018

Support single precision brian-team/brian2genn#69

Closed

mstimberg merged commit b6a3a6c into master Aug 29, 2018

mstimberg deleted the float32_support branch August 29, 2018 12:55

[MRG] Float32 support #981

[MRG] Float32 support #981

Conversation

mstimberg commented Jul 12, 2018

denisalevi commented Aug 9, 2018

mstimberg commented Aug 13, 2018

thesamovar left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

thesamovar commented Aug 24, 2018

mstimberg commented Aug 24, 2018

thesamovar commented Aug 24, 2018

mstimberg commented Aug 24, 2018

mstimberg commented Aug 27, 2018

thesamovar commented Aug 27, 2018

mstimberg commented Aug 28, 2018

tnowotny commented Aug 28, 2018

mstimberg commented Aug 28, 2018

thesamovar commented Aug 28, 2018

mstimberg commented Aug 28, 2018

tnowotny commented Aug 29, 2018

thesamovar left a comment

Choose a reason for hiding this comment

mstimberg commented Aug 29, 2018

mstimberg commented Aug 29, 2018

tnowotny commented Aug 29, 2018 • edited Loading

neworderofjamie commented Aug 29, 2018 • edited Loading

mstimberg commented Aug 29, 2018

neworderofjamie commented Aug 29, 2018

tnowotny commented Aug 29, 2018

tnowotny commented Aug 29, 2018

mstimberg commented Aug 29, 2018

neworderofjamie commented Aug 29, 2018 via email

tnowotny commented Aug 29, 2018

denisalevi commented Aug 29, 2018

tnowotny commented Aug 29, 2018

neworderofjamie commented Aug 29, 2018

neworderofjamie commented Aug 29, 2018

mstimberg commented Aug 29, 2018

mstimberg commented Aug 29, 2018

tnowotny commented Aug 29, 2018 •

edited

Loading

neworderofjamie commented Aug 29, 2018 •

edited

Loading