ENH: add cupy tensor space by kohr-h · Pull Request #1231 · odlgroup/odl

kohr-h · 2017-11-13T22:07:59Z

No description provided.

adler-j · 2017-11-13T23:15:32Z

Looks like a trivial review, I'll try it tomorrow. Does this work with astra data containers?

kohr-h · 2017-11-13T23:34:46Z

Sure, go ahead. It's work in progress though.

kohr-h · 2017-11-13T23:35:13Z

I haven't tried ASTRA interoperability yet.

adler-j

Frankly only some minor stuff left, looks freaking magic to me!

adler-j · 2017-11-14T16:57:57Z

+else:
+    from pkg_resources import parse_version
+    if parse_version(cupy.__version__) < parse_version('2.0.0rc1'):
+        raise ImportError('cupy <2.0.0rc1 not supported')


This should be a warning, we don't want to render ODL unusable due to a wrong version.

The thing is that version 2.0.0rc1 adds support for complex arrays, I really don't want anything below. Actually we should go for 2.0.0, the latest one on PyPI. It wasn't out at the time of writing this.

Okay, complex stuff is available in 2.0.0, I'll emit a warning for earlier versions.

adler-j · 2017-11-14T16:58:33Z

+# --- Space method implementations --- #
+
+
+lico = cupy.ElementwiseKernel(in_params='T a, T x, T b, T y',


Are these lazily evaluated? If not, we should be careful about the startup cost.

They should be, since T is a type template, so the only alternative would be to compile for all types, which I cannot imagine. The good thing is that the kernels are cached on disk, so after the first call things run way faster.

I just checked. When I removed pkg_resources (bloody slow, we should avoid it whenever possible, at least during import odl), the import time of cupy_tensors.py is practically identical to the time of import cupy (~40 ms).

adler-j · 2017-11-14T16:59:49Z

+    elif np.dtype(dtype) == 'float64':
+        prefix = 'd'
+    else:
+        raise ValueError('dtype {!r} not supported by cuBLAS'.format(dtype))


Not even complex values?

cuBLAS itself supports it, maybe the bindings are not implemented yet in cupy. I'll check.

Supported in cupy 2.0.0

adler-j · 2017-11-14T17:01:00Z

+
+class CupyTensorSpace(TensorSpace):
+
+    """Tensor space implemented with GPU arrays.


This should mention what backend is used.

adler-j · 2017-11-14T17:05:08Z

+            self.__weighting = CupyTensorSpaceCustomNorm(norm)
+        elif inner is not None:
+            self.__weighting = CupyTensorSpaceCustomInner(inner)
+        else:  # all None -> no weighing


What happened to our no weighting vs const weighting debate? I.e. in several cases we expect stuff like this to be weighted:

space(weighting=n/10.0) # "randomly" gives no weighting for n=10

I guess here it matters less than for the discretized spaces, but in them it really maters.

spc = odl.uniform_discr(0, 5, n) # 'randomly' not weighted for n=5

this causes problems downstream for e.g. RayTransform which has special behaviour for non-weighted spaces.

I'd still need a convincing argument why weighting with constant 1 would be fundamentally different from no weighting. I currently see it just as an optimization that scraps a multiplication with 1.

Is this not a very weird behaviour:

>>> spc = odl.uniform_discr([0, 0], [9.999, 9.999], [10, 10]) >>> rt = odl.tomo.RayTransform(spc, odl.tomo.parallel_beam_geometry(spc)) >>> rt.range uniform_discr([ 0. , -14.1407], [ 3.1416, 14.1407], (45, 31))

vs

>>> spc = odl.uniform_discr([0, 0], [10, 10], [10, 10]) >>> rt = odl.tomo.RayTransform(spc, odl.tomo.parallel_beam_geometry(spc)) >>> rt.range uniform_discr([ 0. , -14.1421], [ 3.1416, 14.1421], (45, 31), weighting=1.0)

I mean, wat? this has to be considered a rather severe bug, no?

That seems quite wrong.

But part of the problem seems to me that we need to jump through hoops currently to make the adjoint correct for weighted vs. unweighted spaces. Doing it via something like #1177 will solve the issue very elegantly by keeping all weighting stuff completely outside the operator definitions.
And when such a system is in place, no weighting and weighting with 1 will be exactly equivalent.

Well this has nothing to do with the adjoint (in fact we don't even need the adjoint to have this problem), the problem here is that we want "the same weighting scheme" on the range as on the domain, but if we dont distinguish no weighting and 1.0 weight, we can't do that, because there is no concept of "the same weighting".

The solution in this case would be to fully remove the is_weighted flag and remove the "unweighted" functionality for the ray transform, but I'm not sure if I'm happy about that.

Oh right, I was totally forgetting about the fact that we infer the range. Okay now I got it. Hm. We do have an optional range parameter for the ray transform, don't we? I tend to think that the behavior should be

as currently when we have a TensorSpace with cell_volume weighting,

no weighting in any other case.

To make it easier to change this, we could have a asweighted() method on TensorSpace where a new space of the same type but with new weighting can be created.

Well it's not quite obvious to me how we would figure out if we're using cell_volume weighting 1.0 or no weighting unless we expose this flag somehow. We previously did it by exposing thetype of the weighting, but I guess we could fix this by exposing the argument used to create the weighting (somehow), i.e. space.weighting_type returns a string 'const' or w/e. Does that sound reasonable?

I feel adding yet another as_... should hopefully not be needed right now, but I guess in the long run.

Has this been further addressed? Otherwise we need to solve it soon (perhaps not in this PR tho)

adler-j · 2017-11-14T17:10:43Z

+                raise RuntimeError("no conversion for dtype {}"
+                                   "".format(arr.dtype))
+        else:
+            space = type(self.space)(arr.shape, dtype=self.dtype,


I'll port the current implementation in NumpyTensorSpace, the proper solution will come with #1238
Just in general: if you squash axes, there's no obvious way to propagate weightings.

We need to settle what squashed axes mean. Imo it should set that dimension to "1" so to say, i.e. we keep space[0:1].weighting and space[0].weighting different.

Well, as of now, i.e., before #909 is done (part of #1238) we can't know how to handle weighting constants when removing axes. If all axes stay intact, we can keep the constant. If an axis is removed, we have to fall back to default. Not so great, but yeah.
For array weighting, i.e., the weights have the same shape as the space, we can just index them the same way, so that's a bit easier.

adler-j · 2017-11-14T17:11:25Z

+        Numpy implementation is used which causes significant overhead
+        due to data copies between host and device.
+        """
+        # TODO: Test with some native ufuncs, then remove this attribute


fix this todo

This raises a question: If we have a cupy space element and run some numpy code on it that requires a cast to numpy array, should we just go ahead and do it, at the cost of hidden performance loss, or should we raise?
I have a slight preference but want to make sure we're on the same page.

Well as always im in the "at least it works" boat. I guess it comes down to how feature complete cupy is, if it has almost everything it is fine, but if you stumble every thing you do, it's gonna be irritating to work with

I'm in the same boat, so that's fine then. As far as I can see, they're pretty complete on the ufunc front, but methods like at and accumulate are largely not implemented. So we need some workaround to use the CUDA code for the obvious ones (sum, prod, cumsum, cumprod for add and multiply), but otherwise drop out to Numpy.

The cupy folks are a bit more worried about efficiency, probably since in neural nets, things are already buried deep down and hard to debug, so this kind of bottleneck danger would be pretty bad for them.

Go ahead with that then!

I've implemented the manual fiddling.

adler-j · 2017-11-14T17:12:03Z

+        newreal : `array-like` or scalar
+            The new real part for this tensor.
+        """
+        self.real.data[:] = newreal


I think self.data.real[:] would be slightly more efficient?

Very possible.

adler-j · 2017-11-14T17:12:14Z

+        real : `CupyTensor` view with real dtype
+            The real part of this tensor as an element of an `rn` space.
+        """
+        # Only real dtypes currently


Quite sure there were complex ones above?

Needs update obviously.

adler-j · 2017-11-14T17:12:48Z

+    if np.isscalar(weights):
+        weighting = CupyTensorSpaceConstWeighting(weights, exponent=exponent)
+    else:
+        # TODO: sequence of 1D array-likes


Cite the issue number

kohr-h · 2017-11-14T17:27:27Z

+# v. 2.0. If a copy of the MPL was not distributed with this file, You can
+# obtain one at https://mozilla.org/MPL/2.0/.
+
+"""Implementation of tensor spaces using ``pygpu``."""


Obviously the docs need update :-)

kohr-h · 2017-11-14T17:36:05Z

+        >>> same_space == space
+        True
+        """
+        return (super().__eq__(other) and


adler-j · 2017-11-15T08:51:49Z

Ok, so odlcuda is officially broken. I'll try to fix it, but getting this in would be a preferable solution. What is the ETA here, code looks quite good.

kohr-h · 2017-11-15T09:03:44Z

I'll give it a couple of days of focused work, but those days will be distributed a bit. I give it high prio at least.

adler-j · 2017-11-15T09:06:30Z

I'm fixing odlcuda anyway

kohr-h · 2017-11-23T09:56:11Z

I'll look into this today. Also in view of #1246.

kohr-h · 2017-11-23T18:00:24Z

Getting there. I fixed the complex stuff, and all the doctests go green. Unit tests next.

adler-j

Some further comments to take into account. This is looking great!

adler-j · 2017-11-24T12:18:35Z

+    signature_string, indent)
+
+try:
+    import cupy


How long does this import take?

40 ms, it's really fast.

Great! No worries then

adler-j · 2017-11-24T12:41:00Z

+    return lico(a, x, 1, y, y)
+
+
+def _flat_inc(arr):


what does inc mean here?

It's the incx, incy stuff needed for the cuBLAS functions.

adler-j · 2017-11-24T12:41:36Z

+
+
+def _flat_inc(arr):
+    """Compute the flat element stride for cuBLAS if possible, else raise."""


A Parameters, Returns, Raises here would be great

Was working on that this morning. The implementation was also not correct :-)

I redid the cuBLAS handling, works as expected now. Also the lico fallback kernel takes arrays of different dtypes.

adler-j · 2017-11-24T14:04:29Z

+        Numpy implementation is used which causes significant overhead
+        due to data copies between host and device.
+        """
+        # TODO: Test with some native ufuncs, then remove this attribute


Go ahead with that then!

adler-j · 2017-11-24T14:08:52Z

+        rn(3, impl='cupy', weighting=[1, 2, 3])
+        >>> space = odl.tensor_space((2, 3), impl='cupy', dtype=int)
+        >>> space
+        tensor_space((2, 3), 'int', impl='cupy')


Won't these doctests fail misserably for users without cupy, or do we exclude them somehow?

At least not when running the file. I'll have to check how to configure pytest to exclude them as well.

Well if you explicitly run this file, you should get errors IMO, its harder with the global tests

adler-j · 2017-11-24T14:15:28Z

+    @property
+    def data_ptr(self):
+        """A raw pointer to the data container.
+


Returns ------- data_ptr : int

to show the type of the returned value

I prefer not to add Returns to @property docstrings since they do not "feel" like functions that return something. I'll add it somewhere else.

Anything goes as long as the info is there!

adler-j · 2017-11-24T14:19:49Z

+
+        Notes
+        -----
+        The element-by-element comparison is performed on the CPU,


Really? Why not write a ReductionKernel for this? Should be rather simple

Makes sense.

Forgot to update the doc?

Edit: that was indeed the case. I fixed the doc

adler-j · 2017-11-24T14:22:35Z

+        >>> x[::2]
+        rn(3, impl='cupy').element([ 1.,  3.,  5.])
+
+        The returned views are writable, so modificatons alter the


Also mention that this is not always the case (i guess?)

Not for index arrays and boolean arrays, right.

adler-j · 2017-11-24T14:24:04Z

+                raise RuntimeError("no conversion for dtype {}"
+                                   "".format(arr.dtype))
+        else:
+            space = type(self.space)(arr.shape, dtype=self.dtype,


We need to settle what squashed axes mean. Imo it should set that dimension to "1" so to say, i.e. we keep space[0:1].weighting and space[0].weighting different.

adler-j · 2017-11-24T14:26:37Z

+
+
+if CUPY_AVAILABLE:
+    dotw = cupy.ReductionKernel(in_params='T x, T y, W w',


Are these compiled on the fly later? We don't want to create startup latency

Yes, as I wrote somewhere else, the import time of this module is practically identical to the import time of cupy, around 40 ms. So these guys don't add overhead at creation time.

…,accumulate)

kohr-h · 2017-12-01T08:52:36Z

-        tmp = np.empty(out.shape, out.dtype, order=out.space.default_order)
-        with writable_array(out) as out_arr:
+        tmp = array_module(impl).empty(
+            out.shape, out.dtype, order=out.space.default_order)


This of course assumes that impl has the same interface as numpy.

Perhaps it's better to implement wrappers for the most common functions, like asarray (already done), empty, and whatever else is needed. That would be more extensible.

That would be how e.g. Keras does it, but I say we leave this for now. Changing should be rather easy (just search replace basically).

kohr-h · 2017-12-01T08:55:49Z

+            Array into which the result should be written. Must be contiguous
+            and of the correct dtype.
+        impl : str, optional
+            Array backend for the output, used when ``out`` is not given.


Should probably mention where the valid choices come from.

kohr-h · 2017-12-01T08:56:35Z


        pspace = self
-        for n in shape:
+        for n in reversed(shape):


This was an actual bug.

kohr-h · 2017-12-01T08:58:07Z

-    def asarray(self, out=None):
-        """Extract the data of this array as a ``numpy.ndarray``.
+    def asarray(self, out=None, impl='numpy'):
+        """Extract the data of this array as an ndarray.


Perhaps make a glossary term ndarray that explains the different flavors.

kohr-h · 2017-12-01T09:10:28Z

            return float(np.sqrt(self.const) * _norm_default(x))
        elif self.exponent == float('inf'):
-            return float(self.const * _pnorm_default(x, self.exponent))
+            return float(_pnorm_default(x, self.exponent))


I've changed this because there are two reasons for it:

The mathematical one: ||x||_{p, w} --> ||x||_{inf, w} should hold for p --> inf (we knew that).

Also structure-wise, all unweighted p-norms follow the pattern

norm = post_map( reduce( map(x) ) )

e.g.

1 <= p < inf: map(x) = abs(x) ** p, reduce(x) = sum(x), post_map(x) = x ** (1/p)

p = inf: map(x) = abs(x), reduce(x) = max(x), post_map(x) = x

Weighted norms have the structure

norm = post_map( reduce( weight * map(x) ) )

and here it follows naturally that the weights play no role for p = inf since they do not change the max. Why treat this case completely differently?

I guess you're right, lets do it like this for now.

kohr-h · 2017-12-01T09:12:39Z

    def is_valid(self):
        """Return True if the array is a valid weight, i.e. positive."""
-        return np.all(np.greater(self.array, 0))
+        return (self.array > 0).all()


Again assuming a certain API of the arrays. Perhaps better to have wrapper all, any and such. Dunno.

Lets leave it like this until someone wants to add something that breaks it

kohr-h · 2017-12-01T09:19:51Z

+for name, nin, nout, docstring in UFUNCS:
+    if nin == 2:
+        # Currently not supported
+        continue


We had a bunch of dysfunctional ufunc ops created here.

kohr-h · 2017-12-01T09:21:03Z

+        if isinstance(iter2, CupyTensor):
+            iter2 = iter2.asarray()
+        elif isinstance(iter2, cupy.ndarray):
+            iter2 = cupy.asnumpy(iter2)


Should use utility.asarray here I guess.

kohr-h · 2017-12-01T09:22:11Z



+def xfail_if(condition, reason=''):
+    """Return a ``pytest.xfail`` object if ``condition`` is ``True``."""


Explain usage.

adler-j

Massive review. Mostly looking good, just some minor stuff.

With that said, the increase in test run-time is worrying. I guess we can live with it for now, but we must diagnoize and improve the situation.

adler-j · 2017-12-01T00:38:10Z


        pspace = self
-        for n in shape:
+        for n in reversed(shape):


Why this? Wouldn't our users expect r2 ** (3, 4) == (r2 ** 3) ** 4?

I don't think so. When doing space[i], the index slices along the outermost power, i.e., for space ** (3, 4) you expect the valid indices i to be 0, 1, 2, no? It's the same logic as (R^m)^n == R^(n x m).
In your example, I would expect that r2 ** (3, 4) has shape (3, 4, 2).

I concur with your statement. It certainly makes sense.

adler-j · 2017-12-01T00:39:54Z

+    be applied, a ``ValueError`` is raised, triggering a fallback
+    implementation.
+
+    For **1 array**, the conditions to be fulfilled are


adler-j · 2017-12-01T00:42:17Z

+        return _cublas_scal(
+            x.data.device.cublas_handle, x.size, a, x.data.ptr, incx)
+
+    scal.__name__ = scal.__qualname__ = '_cublas_scal'


Why not simply name it _cublas_scal?

Renamed this.

adler-j · 2017-12-01T00:42:50Z

+
+    If possible, a cuBLAS implementation is returned, otherwise a fallback.
+
+    In general, cuBLAS requires single or double precision float or complex


This string is overly long for these, simply reference the function above.

adler-j · 2017-12-01T00:44:07Z

+    This implementation is a highly optimized, considering all special
+    cases of array alignment and special scalar values 0 and 1 separately.
+    """
+    scal1 = _get_scal(x1.data)


Perhaps we should consider some type of cache for these? Or is that implemented downstream?

Cache is an optimization that will have to come later. I'll focus on getting this in first.

Agreed, we have the same thing on the Numpy spaces, we can add caches for both in one go.

adler-j · 2017-12-05T00:05:30Z

+for name, nin, nout, docstring in UFUNCS:
+    if nin == 2:
+        # Currently not supported
+        continue


adler-j · 2017-12-05T00:06:00Z

    globals()[name] = ufunc_factory
    __all__ += (name,)

+np.seterr(**npy_err_old)


We should honestly add some "we're so sorry" line above this :D

adler-j · 2017-12-05T00:07:11Z

+        if isinstance(iter2, CupyTensor):
+            iter2 = iter2.asarray()
+        elif isinstance(iter2, cupy.ndarray):
+            iter2 = cupy.asnumpy(iter2)


adler-j · 2017-12-05T00:08:45Z

+    from odl.space.pspace import ProductSpace
+
+    if isinstance(space, ProductSpace) and not space.is_power_space:
+        raise ValueError('`space` cannot be a non-power product space')


Why not? Part of the reason we use this function is that it supports these things (as compared to e.g. space.element(np.random.rand))

adler-j · 2017-12-05T00:12:34Z

+    # Workaround for `shape` not using the base space shape of a
+    # power space
+    # TODO: remove when fixed, see
+    # https://github.com/odlgroup/odl/pull/1152


I'll get this in ASAP

This is now in, yay!

adler-j · 2018-02-07T10:56:11Z

So, I'm officially announcing that I'll "take over" this PR and make a new PR for it.

kohr-h · 2018-02-07T12:48:00Z

So, I'm officially announcing that I'll "take over" this PR and make a new PR for it.

Go for it! Maybe close this one for difficulty to open the page?

kohr-h · 2018-02-07T12:56:03Z

+    ValueError
+
+    Slicing in the **fastest** axis is allowed as it results in a constant
+    stride in the flat memory. Slicing with stride in any other axis is


Oops, "is" shouldn't be there

adler-j · 2018-02-07T13:14:42Z

I'll close this once I've fixed the outstanding issues

mehrhardt · 2018-04-12T10:45:54Z

Any news on this front?

n1kt0 · 2019-05-16T14:20:22Z

How far is the integration process?

Thanks in advance and best regards,

Nikita

adler-j · 2019-05-16T15:13:18Z

This has sadly stalled quite badly and I will not have time to focus on it in the coming months. If anyway wants to pick up I'd be more than happy.

n1kt0 · 2019-05-16T15:24:39Z

Are there some tests or do you have a summary what needs to be done or where one could start? Am Do., 16. Mai 2019 um 17:13 Uhr schrieb Jonas Adler < notifications@github.com>:

…

This has sadly stalled quite badly and I will not have time to focus on it in the coming months. If anyway wants to pick up I'd be more than happy. — You are receiving this because you commented. Reply to this email directly, view it on GitHub <#1231?email_source=notifications&email_token=AAJ3R3JJGRZ526MMSBPVQD3PVV2Y7A5CNFSM4EDR6TCKYY3PNVWWK3TUL52HS4DFVREXG43VMVBW63LNMVXHJKTDN5WW2ZLOORPWSZGODVSEDPQ#issuecomment-493109694>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AAJ3R3OJQDNXRJL3K2G4WYTPVV2Y7ANCNFSM4EDR6TCA> .

adler-j · 2019-05-20T09:42:50Z

Sadly there is no summary but the main issue is to get the tests running. With that said this PR has become so old that it might be worth re-starting it with the newly updated spaces, perhaps @kohr-h knows how much has changes that touches upon this?

kohr-h · 2019-05-20T10:00:54Z

I'm hesitating to re-activate this one before we have settled on #1459 and #1458. Depending on the outcome of the latter in particular, getting in support for CuPy will be massively simpler (or not).

adler-j mentioned this pull request Nov 14, 2017

Release 1.0.0 #1055

Open

20 tasks

adler-j reviewed Nov 14, 2017

View reviewed changes

kohr-h commented Nov 14, 2017

View reviewed changes

adler-j mentioned this pull request Nov 22, 2017

tensor update somehow breaks with cuda #1246

Open

Holger Kohr added 3 commits November 23, 2017 15:01

ENH: add cupy tensor space

d028dd5

MAINT: improve import of cupy

dd2b7c4

WIP: fix cupy->numpy transfers

e1ce7fb

kohr-h force-pushed the cupy_space branch from 4fc8c3b to e1ce7fb Compare November 23, 2017 17:59

Holger Kohr added 3 commits November 24, 2017 00:14

WIP: fix custom cupy kernels

f282947

MAINT: minor fixes

b84ee8d

WIP: implement cupy space unit tests

4a0c67c

adler-j reviewed Nov 24, 2017

View reviewed changes

Holger Kohr added 9 commits November 24, 2017 15:36

WIP: fix cublas stuff

1d431c6

WIP: fix cublas, weight propagation, implement (add,multiply).(reduce…

2fdebff

…,accumulate)

TST: ignore cupy_tensors in doctests if cupy is not available

3bdc3f0

MAINT: fix various numpy and cupy tensor things

6fbcc81

ENH: allow indexing of CupyTensor with CupyTensor

f6a8022

MAINT: minor stuff

ee386ab

MAINT: change weighted inf-norm to ignore weight

2b7103c

ENH: make comparison of cupy arrays in tests faster

2ed9ae1

WIP: fix tensor tests for cupy

8414458

kohr-h changed the title ~~WIP: add cupy tensor space~~ ENH: add cupy tensor space Nov 30, 2017

TST: mark cupy-only test as skip if cupy is not available

58b921d

kohr-h commented Dec 1, 2017

View reviewed changes

Comment thread odl/set/space.py

pspace = self

for n in shape:

for n in reversed(shape):

kohr-h Dec 1, 2017

Copy link
Copy Markdown

Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This was an actual bug.

kohr-h commented Dec 1, 2017

View reviewed changes

MAINT: fix Py2-incompatible doctest in utility

eb10834

kohr-h mentioned this pull request Dec 1, 2017

ENH: Tensor-valued DiscreteLp #1238

Open

21 tasks

adler-j requested changes Dec 5, 2017

View reviewed changes

kohr-h commented Feb 7, 2018

View reviewed changes

kohr-h removed the status: review needed label Feb 14, 2018

adler-j mentioned this pull request Jun 28, 2018

Add Cupy space #1401

Open

odlgroup deleted a comment from n1kt0 May 16, 2019

kohr-h mentioned this pull request Mar 30, 2020

Configure GPU index for 'astra_cuda', select GPU currently used by PyTorch in OperatorModule #1546

Open

kohr-h mentioned this pull request May 1, 2020

Pytorch and tensor flow backend pass through the CPU. #1558

Open

		# --- Space method implementations --- #


		lico = cupy.ElementwiseKernel(in_params='T a, T x, T b, T y',


		class CupyTensorSpace(TensorSpace):

		"""Tensor space implemented with GPU arrays.



		def _flat_inc(arr):
		"""Compute the flat element stride for cuBLAS if possible, else raise."""



		if CUPY_AVAILABLE:
		dotw = cupy.ReductionKernel(in_params='T x, T y, W w',



		def xfail_if(condition, reason=''):
		"""Return a ``pytest.xfail`` object if ``condition`` is ``True``."""


		If possible, a cuBLAS implementation is returned, otherwise a fallback.

		In general, cuBLAS requires single or double precision float or complex

Conversation

kohr-h commented Nov 13, 2017

Uh oh!

adler-j commented Nov 13, 2017

Uh oh!

kohr-h commented Nov 13, 2017

Uh oh!

kohr-h commented Nov 13, 2017

Uh oh!

adler-j 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!

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!

Choose a reason for hiding this comment

Uh oh!

adler-j Nov 14, 2017 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

adler-j Nov 14, 2017 • edited Loading Uh oh! There was an error while loading. Please reload this page.

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!

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!

kohr-h Nov 23, 2017 • edited Loading Uh oh! There was an error while loading. Please reload this page.

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!

Choose a reason for hiding this comment

Uh oh!

adler-j Nov 14, 2017 •

edited

Loading

adler-j Nov 14, 2017 •

edited

Loading

kohr-h Nov 23, 2017 •

edited

Loading