Add half precision base support #3439

crtrott · 2020-09-30T16:53:43Z

If possible we should move that in after the backend refactor since this file might be one of the per backend specializations.

crtrott · 2020-09-30T16:58:27Z

Forgot the license notice ...

DavidPoliakoff · 2020-09-30T17:01:49Z

core/src/Kokkos_Half.hpp

+#include <cstdint>
+#include <Cuda/Kokkos_Cuda_Half.hpp>
+
+#ifndef KOKKOS_IMPL_HALF_TYPE_DEFINED


Should this be IMPL? I feel like users might want to know whether their Kokkos has half-precision (though I'm open to other mechanisms)

half_t will always be defined, the way to figure out whether we use the fallback is half_is_float constexpr bool.

@crtrott and I discussed in Slack. I was worried about a model where a code used different versions of Kokkos (including one before we make this change), and they might switch off functionality based on that macro. We're going with a model where codes use a version of Kokkos that evolves with the code, so this won't come up

core/src/Cuda/Kokkos_Cuda_Half.hpp

calewis · 2020-09-30T18:30:34Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+KOKKOS_INLINE_FUNCTION
+half_t cast_to_half(float val) { return __float2half(val); }
+KOKKOS_INLINE_FUNCTION
+half_t cast_to_half(double val) {


Do we care that double2half(val) and float2half(static_cast<float>(val)) don't have to return identical values (at least this is true for arm-gcc)?

Isn't that expected?

Users might be surprised if on one platform they get different results from cast_to_half than on another.

that is the nature of half precision right now. To the degree vendors agree on conversions we should be good. But CUDA doesn't provide us the conversion calls on both sides so I rather call the fast one on the device and do the extra conversion (with potentially different result) on host. I think you argument above though that is just generally expected. static_cast<int>(val) doesn't necessarily give the same result as static_cast<int>(static_cast<float>(val)) either.

core/src/Cuda/Kokkos_Cuda_Half.hpp

dalg24 · 2020-09-30T19:54:43Z

core/unit_test/TestHalfConversion.hpp

+
+template <class T>
+void test_half_conversion_type() {
+  double epsilon = Kokkos::Experimental::half_is_float ? 0.0000003 : 0.0003;


How did you come up with this relative tolerance?

wikipedia :-) This is the actual (or pretty close) value based on the number of bits for the mantisse.

dalg24 · 2020-09-30T21:14:46Z

core/unit_test/TestHalfConversion.hpp

+  T base         = static_cast<T>(3.3);
+  Kokkos::Experimental::half_t a = Kokkos::Experimental::cast_to_half(base);
+  T b                            = Kokkos::Experimental::cast_from_half<T>(a);
+  ASSERT_TRUE((double(b - base) / double(base)) < epsilon);


Did you look into gtest support for fp comparison?

Use ASSERT_NEAR

hm I am not sure that that works, since I want to explicitly cast b-base not b and base and then do the difference.

dalg24 · 2020-10-01T21:58:39Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+#include <Kokkos_Macros.hpp>
+#ifdef KOKKOS_ENABLE_CUDA
+#include <cuda_fp16.h>
+#include <cstdint>


I don't think you need it any more

Suggested change

#include <cstdint>

dalg24 · 2020-10-01T22:05:27Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+KOKKOS_INLINE_FUNCTION
+half_t cast_to_half(double val) {
+  // double2half was only introduced in CUDA 11 too
+  return __float2half(static_cast<float>(val));


I would prefer if you wrote [unsigned] short int, [unsigned] long long int

(not omitting the trailing "int")

dalg24 · 2020-10-01T22:10:34Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+template <class T>
+KOKKOS_INLINE_FUNCTION
+    typename std::enable_if<std::is_same<T, float>::value, T>::type


Suggested change

typename std::enable_if<std::is_same<T, float>::value, T>::type

std::enable_if_t<std::is_same<T, float>::value, T>

ah right I am forgetting that we now got C++14

lets change all of those.

core/src/Kokkos_Half.hpp

crtrott · 2020-10-05T20:28:56Z

There are more changes coming to this. In particular for the CUDA version we need our own struct which does the right thing on the host and the device (i.e. overloaded math operators). Later we will add more versions: Intel only, CUDA on ARM (where we can actually use __fp16 on the host), HIP, etc.

Adds conversion functions.

Adds operator overloads for half_t. Adds support for half_t arithmetic on host via promotion to float.

Add Cuda half_t type and operators

crtrott

There are a number of changes which need to be done see my comments.

crtrott · 2020-10-13T15:03:38Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  operator half_device_type() const { return val; }
+
+// NOTE: Changing below to 1 produces constructor overload error
+#if 0


remove this.

ok we need to figure out why this is still an issue.

crtrott · 2020-10-13T16:01:22Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+namespace Kokkos {
+namespace Experimental {
+
+using half_device_type = __half;


I think this should be an internal typedef to half_t and it should be called impl_type.

crtrott · 2020-10-13T16:02:44Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+  KOKKOS_FUNCTION
+  half_t(half_device_type rhs = 0) : val(rhs) {}
+


maybe the bool thing gets solved if we explicitly add half_t(const half_t& rhs):val(rhs.val) {} constructor?

crtrott · 2020-10-13T16:03:56Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  half_t operator+() const {
+    half_t tmp = *this;
+#ifdef __CUDA_ARCH__
+    // printf("half_t unary operator+\n");


remove all the printf statements.

crtrott · 2020-10-13T16:26:54Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+  // Logical operators
+  KOKKOS_FUNCTION
+  half_t operator!() const {


this needs to return bool. Which also means we don't need the tmp but simply return !val or !__half2float(val)

crtrott · 2020-10-13T16:27:19Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+    return tmp;
+  }
+
+#if 1


get rid of the #if 1

crtrott · 2020-10-13T16:27:31Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  // NOTE: Loses short-circuit evaluation
+  KOKKOS_FUNCTION
+  bool operator&&(half_t rhs) const {
+    half_t tmp = *this;


don't need a tmp here.

crtrott · 2020-10-13T16:27:44Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  // NOTE: Loses short-circuit evaluation
+  KOKKOS_FUNCTION
+  bool operator||(half_t rhs) const {
+    half_t tmp = *this;


don't need a tmp here.

crtrott · 2020-10-13T16:27:54Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  // Comparison operators
+  KOKKOS_FUNCTION
+  bool operator==(half_t rhs) const {
+    half_t tmp = *this;


don't need a tmp here.

crtrott · 2020-10-13T16:28:49Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+template <class T>
+KOKKOS_INLINE_FUNCTION
+    typename std::enable_if<std::is_same<T, float>::value, T>::type


lets change all of those.

crtrott · 2020-10-14T15:17:04Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  }
+
+  template <class T>
+  KOKKOS_FUNCTION half_t operator=(T rhs) {


should return reference.

e10harvey · 2020-10-14T16:55:14Z

I've attempted to address the CI build errors and feedback in crtrott#7.

Implement PR feedback

nliber · 2020-10-15T00:19:47Z

core/src/Kokkos_Half.hpp

+// Using an explicit list here too, since the other ones are explicit and for
+// example don't include char
+template <class T>
+KOKKOS_INLINE_FUNCTION typename std::enable_if<


Minor: we can use enable_if_t here (C++14).

nliber · 2020-10-15T00:27:56Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  half_t(const half_t&) = default;
+
+  KOKKOS_FUNCTION
+  half_t(impl_type rhs = cast_to_half(0)) : val(rhs) {}


I have a slight preference for two constructors here instead of relying on defaulted parameters.

nliber · 2020-10-15T00:30:09Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+  // Cast rhs to half for assignment to lhs of type half_t
+  template <class T>
+  KOKKOS_FUNCTION half_t(T rhs) : half_t(cast_to_half(rhs)) {}


I'm wondering if we should constrain this (otherwise it matches just about everything) instead of the hard error calling it on something not convertible to a parameter that cast_to_half can take.

nliber · 2020-10-15T00:37:17Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  }
+
+  KOKKOS_FUNCTION
+  half_t(const half_t&) = default;


I'm thinking Rule of 0 here. I can't imagine move does anything different vs. copy for __half.

we had some issues with ambiguity before anything to consider in that direction? Or are you just saying we shouldn't define any of those constructors, and just write the explicit conversion constructors explicitly.

nliber · 2020-10-15T02:11:32Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+  // Binary Arithmetic
+  KOKKOS_FUNCTION
+  half_t operator+(half_t rhs) const {


These binary operators should be hidden friends, as in:

half_t friend operator+(half_t lhs, half_t rhs) { #ifdef __CUDA_ARCH__ lhs.val += rhs.val; #else lhs.val = __float2half(__half2float(lhs.val) + __half2float(rhs.val)); #endif return lhs; }

The benefits are symmetry and ADL.

It is less common to do those for unary operators, although Anthony Williams makes a case for it at https://www.justsoftwaresolutions.co.uk/cplusplus/hidden-friends.html.

Remove implicit conversion from half_t to __half and bool. This may cause the compiler to consider all __half operators or bool operators for expressions involving half_t. Make operator{+,-,*,/} symmetric Explicitly overload half_t constructors Add casting to/from bool Add test cases for symmetry Do not force a copy constructor for half_t

Implement PR feedback and implicit ops

Add explicit conversion ops for casting from half_t to supported types Add copy constructor back to make half_t both trivially and copy constructible

Use explicit casting to support deep_copy from T to half_t

half_t Cuda and deep_copy updates

CI workarounds and fixes

Add back type_traits include...

apply-clang-format...

Try KOKKOS_ENABLE_SYCL to disable the lambda...

Conditionally disable TestHalfOperators

Test whether half_t is trivially copyable

dalg24

Drive-by review. I need more time.

dalg24 · 2020-10-13T16:28:13Z

core/unit_test/TestHalfConversion.hpp

+  T base         = static_cast<T>(3.3);
+  Kokkos::Experimental::half_t a = Kokkos::Experimental::cast_to_half(base);
+  T b                            = Kokkos::Experimental::cast_from_half<T>(a);
+  ASSERT_TRUE((double(b - base) / double(base)) < epsilon);


Use ASSERT_NEAR

dalg24 · 2020-10-23T02:51:48Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, float>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, bool>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, double>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, short>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, int>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, long>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, long long>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION
+    std::enable_if_t<std::is_same<T, unsigned short>::value, T>
+        cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_same<T, unsigned int>::value, T>
+    cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION
+    std::enable_if_t<std::is_same<T, unsigned long>::value, T>
+        cast_from_half(half_t);
+template <class T>
+KOKKOS_INLINE_FUNCTION
+    std::enable_if_t<std::is_same<T, unsigned long long>::value, T>
+        cast_from_half(half_t);


Readability is poor. Why did you change from

std::enable_if_t<std::is_same<T, Foo>::value && std::is_same<T, Bar>::value && ... , T>

Good point. I will update accordingly. I used a regex to convert the cast_to_half forward decls above.

@dalg24: Thinking about this more now. I cannot overload based on return type only. This was the best way I could find to provide the forward decls for cast_from_half. Is there a better way to do this?

std::enable_if_t<std::is_same<T, Foo>::value && std::is_same<T, Bar>::value && ... , T>

I don't think this expression will ever evaluate to true.

should be && -> ||

I get a multiple definition error with ||. Is there a better way to provide the forward decls than what is shown above?

dalg24 · 2020-10-23T02:53:06Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+namespace Kokkos {
+namespace Experimental {
+#define HALF_IMPL_TYPE __half


Why a macro?

I think there is a better way to do this. I needed a compile time conditional since using impl_type = T was not possible in Kokkos_Half.hpp. Any suggestions?

I am not quite following why using a type alias is not possible.
In any case you would probably need to comment to that effect if you absolutely have to use a macro and also make sure you #undef it when you don't need it any more.

I have a suggestion, this is mostly about the test thing there we want to test interoperability with implementation types:
We can simply have a trait to do that however:

template<class T> struct half_impl_type { using type = typename T::impl_type; }; template<> struct half_impl_type<float> { using type = float; };

I would put that into the Impl namespace for now.

dalg24 · 2020-10-23T02:56:53Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+  KOKKOS_FUNCTION
+  half_t(impl_type rhs) : val(rhs) {}
+  KOKKOS_FUNCTION
+  half_t(float rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(bool rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(double rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(short rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(int rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(long rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(long long rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(unsigned short rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(unsigned int rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(unsigned long rhs) : val(cast_to_half(rhs).val) {}
+  KOKKOS_FUNCTION
+  half_t(unsigned long long rhs) : val(cast_to_half(rhs).val) {}


Why are these not explicit?

This is a best faith effort to get half_t behaving like other precisions (ex: float). I observed that float can be implicitly cast to and from T. For half_t, we cannot have implicit conversion operators to T as the compiler will consider T and its operators in all expressions involving half_t; but, we cannot mark half_t's operators as higher precedence than T's.

Fix half_t trivially copyable test

e10harvey · 2020-10-27T15:40:02Z

@dalg24, @nliber: Thank you for your reviews.

@dalg24: Do you have any other feedback before I start addressing it?

@nliber: I've attempted to address your feedback in
15104e6. Is there anything else you'd like addressed?

Implement PR feedback

nliber

Looks good, other than the question about short circuiting operators.

nliber · 2020-10-15T02:20:55Z

core/src/Cuda/Kokkos_Cuda_Half.hpp

+
+  // NOTE: Loses short-circuit evaluation
+  KOKKOS_FUNCTION
+  bool operator&&(half_t rhs) const {


Do we really want this? Losing short circuiting can lead to bugs in complex expressions.

Thank you, @nilber. The unit tests were written to check for simple short circuiting. I will update the unit tests with some more complex expressions that rely on operator precedence as well.

what kind of bugs? We could collect feedback on this. But since we don't have implicit conversion to bool (only from) statements like a&&b wouldn't work otherwise and you would need to do bool(a)&&bool(b) instead. So I am in favor of this.

In favor of keeping these logical operator overloads? Shall I update the unit tests with some more complex expressions involving these logical operators?

I think I just merge it. Nevin did approve after all. And if its a real issue he shouldn't have :-)

e10harvey · 2020-11-02T14:32:48Z

@crtrott, @dalg24, @nliber: Is there anything else that should be addressed here before we squash and merge?

crtrott force-pushed the add-half-precision-base-support branch from 578c4e4 to abbfe23 Compare September 30, 2020 16:58

crtrott mentioned this pull request Sep 30, 2020

Initial support for half precision kokkos/kokkos-kernels#794

Merged

DavidPoliakoff reviewed Sep 30, 2020

View reviewed changes

crtrott force-pushed the add-half-precision-base-support branch 2 times, most recently from 7d178f8 to cfe944f Compare September 30, 2020 17:28

dalg24 reviewed Sep 30, 2020

View reviewed changes

core/src/Cuda/Kokkos_Cuda_Half.hpp Outdated Show resolved Hide resolved

core/src/Cuda/Kokkos_Cuda_Half.hpp Outdated Show resolved Hide resolved

core/src/Cuda/Kokkos_Cuda_Half.hpp Outdated Show resolved Hide resolved

calewis reviewed Sep 30, 2020

View reviewed changes

crtrott force-pushed the add-half-precision-base-support branch from cfe944f to 5339250 Compare September 30, 2020 19:13

dalg24 reviewed Sep 30, 2020

View reviewed changes

core/src/Cuda/Kokkos_Cuda_Half.hpp Show resolved Hide resolved

dalg24 reviewed Sep 30, 2020

View reviewed changes

crtrott force-pushed the add-half-precision-base-support branch from 5339250 to 288ab36 Compare October 1, 2020 21:18

dalg24 reviewed Oct 1, 2020

View reviewed changes

crtrott added 2 commits October 5, 2020 13:37

Adding Half precision type for CUDA and float fallback

5cb888d

Adds conversion functions.

Adding test for Half Precision conversions

e7fae55

crtrott force-pushed the add-half-precision-base-support branch from 288ab36 to e7fae55 Compare October 5, 2020 20:38

e10harvey and others added 3 commits October 13, 2020 07:01

Update Half Precision support for CUDA

6873f96

Adds operator overloads for half_t. Adds support for half_t arithmetic on host via promotion to float.

Adds Half Precision operator tests

c0e0372

Merge pull request #6 from e10harvey/add-half-precision-base-support

850b232

Add Cuda half_t type and operators

crtrott commented Oct 13, 2020

View reviewed changes

crtrott commented Oct 14, 2020

View reviewed changes

e10harvey added 2 commits October 14, 2020 09:29

Implement PR feedback

71c3281

Fix travis CI builds

0b82aae

Merge pull request #7 from e10harvey/add-half-precision-base-support

956e9a7

Implement PR feedback

nliber reviewed Oct 15, 2020

View reviewed changes

e10harvey and others added 2 commits October 15, 2020 13:43

Merge pull request #8 from e10harvey/add-half-precision-base-support

b7869ba

Implement PR feedback and implicit ops

e10harvey and others added 15 commits October 16, 2020 15:39

Half precision updates on Cuda

f32d184

Add explicit conversion ops for casting from half_t to supported types Add copy constructor back to make half_t both trivially and copy constructible

Update deep_copy use explicit cast

a369471

Use explicit casting to support deep_copy from T to half_t

Merge pull request #9 from e10harvey/add-half-precision-base-support

d25cb67

half_t Cuda and deep_copy updates

CI workarounds and fixes

43bb63d

Merge pull request #10 from e10harvey/add-half-precision-base-support

aa628e2

CI workarounds and fixes

Add back type_traits include...

9377eb2

Merge pull request #11 from e10harvey/add-half-precision-base-support

072a8bc

Add back type_traits include...

apply-clang-format...

8d9a847

Merge pull request #12 from e10harvey/add-half-precision-base-support

1904454

apply-clang-format...

Try KOKKOS_ENABLE_SYCL to disable the lambda...

ee229b3

Merge pull request #13 from e10harvey/add-half-precision-base-support

3dcbc4a

Try KOKKOS_ENABLE_SYCL to disable the lambda...

Conditionally disable TestHalfOperators

e23c1c5

Merge pull request #15 from e10harvey/add-half-precision-base-support

5967ced

Conditionally disable TestHalfOperators

Test whether half_t is trivially copyable

51d0e36

Merge pull request #16 from e10harvey/add-half-precision-base-support

c8ec1d5

Test whether half_t is trivially copyable

dalg24 reviewed Oct 23, 2020

View reviewed changes

e10harvey and others added 2 commits October 23, 2020 05:20

Fix half_t trivially copyable test

4a0eaa6

Merge pull request #17 from e10harvey/add-half-precision-base-support

eb56dae

Fix half_t trivially copyable test

e10harvey and others added 4 commits October 27, 2020 13:41

Implement PR feedback

0b8eb28

Implement PR feedback

07d5136

Define macro if Cuda has half_t

9a875b4

Merge pull request #18 from e10harvey/add-half-precision-base-support

d80e7c6

Implement PR feedback

nliber approved these changes Oct 28, 2020

View reviewed changes

e10harvey mentioned this pull request Nov 2, 2020

Require clang version 9 or higher when building Kokkos? #3549

Closed

crtrott merged commit f43169d into kokkos:develop Nov 2, 2020

This was referenced Nov 3, 2020

clang/8 + cuda/10.0 build failures following CUDA graph merge #3559

Closed

cuda/9.2 compilation errors: more than one conversion function from "Kokkos::Experimental::half_t::impl_type" to a built-in type #3561

Closed

	typename std::enable_if<std::is_same<T, float>::value, T>::type
	std::enable_if_t<std::is_same<T, float>::value, T>


		KOKKOS_FUNCTION
		half_t(half_device_type rhs = 0) : val(rhs) {}

Add half precision base support #3439

Add half precision base support #3439

Conversation

crtrott commented Sep 30, 2020

crtrott commented Sep 30, 2020

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

crtrott Oct 1, 2020 • edited

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

dalg24 Oct 1, 2020 • edited

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

crtrott commented Oct 5, 2020

crtrott 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

Choose a reason for hiding this comment

e10harvey commented Oct 14, 2020

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

dalg24 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

Choose a reason for hiding this comment

Choose a reason for hiding this comment

e10harvey commented Oct 27, 2020

nliber left a comment • edited

Choose a reason for hiding this comment

nliber Oct 15, 2020 • edited

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

e10harvey commented Nov 2, 2020

crtrott Oct 1, 2020 •

edited

dalg24 Oct 1, 2020 •

edited

nliber left a comment •

edited

nliber Oct 15, 2020 •

edited