Address guvectorize too slow for cuda target by testhound · Pull Request #6889 · numba/numba

testhound · 2021-04-02T04:29:16Z

This pull request address: #6159. It addresses the issues by detecting when the blocksize for a kernel would be < 2 * Number of SM's. I also added a environment variable called NUMBA_CUDA_LOW_OCCUPANCY_WARNINGS that is enabled by default. I had to modify the the class CUDATestCase to disable the warning as many of the test cases trigger this warning.

sklam · 2021-04-02T14:36:43Z

Started smoketest: numba_smoketest_cuda_148

EDIT: should be numba_smoketest_cuda_yaml_36. used a wrong DAG

gmarkall

This is looking promising; there are some comments on the diff suggesting some changes.

It's good that there is a test ensuring this fires for inefficient uses of @guvectorize functions. I think there should also be a test of this using a normal kernel launch - I can't see a really good place to put this new test, but either in numba/cuda/tests/cudapy/test_errors.py or in a new file test_warnings.py.

gmarkall · 2021-04-06T09:49:11Z


+.. envvar:: NUMBA_CUDA_LOW_OCCUPANCY_WARNINGS
+
+   Enable warnings if the block size is too small relative to the number of


As I understand it, the warning is generated when the grid size is too small, not the block size. It's probably worth mentioning that the heuristic used here is (grid size) < 2 * (number of SMs), and that the absence of the warning doesn't necessarily imply a good grid size - larger multiples of the SM count may be more efficient.

It may also be worth mentioning that disabling the warning will reduce the number of CUDA API calls, as it needs to check the number of SMs for the device in the current context to work out if occupancy is low (see comments below on https://github.com/numba/numba/pull/6889/files#diff-b44aea5e6ce50ecba71c563dda6f60805b174c5c048d9287c056d0f8d506a27fR790).

Also it should be mentioned that this is on by default.

gmarkall · 2021-04-06T09:50:46Z

+        # whether to warn about kernels where the block size will
+        # result in a inefficient mapping to GPUs. On by default.


Perhaps "Whether to warn about kernel launches where the grid size will underutilize the GPU due to low occupancy. On by default."?

gmarkall · 2021-04-06T10:07:23Z

+        num_blocks = griddim[0] * griddim[1] * griddim[2]
+        ctx = get_context()
+        smcount = ctx.device.MULTIPROCESSOR_COUNT


I think I'd push this calculation inside an if config.LOW_OCCUPANCY_WARNINGS: condition - if a user is trying to reduce the number of CUDA API calls (get_context can be a little heavy in some use cases) they could then turn off LOW_OCCUPANCY_WARNINGS to save the call to get_context().

gmarkall · 2021-04-06T10:51:17Z

    """
+    def setUp(self):
+        # Disable warnings about low gpu utilization for test suite
+        config.LOW_OCCUPANCY_WARNINGS = 0


This permanently disables low occupancy warnings. To use a setUp() method for disabling the warning for CUDA test cases the following additions are needed:

Memoizing the previous value of config.LOW_OCCUPANCY_WARNINGS in the setUp() method.

A tearDown() method that restores the old value of config.LOW_OCCUPANCY_WARNINGS after the test has run.

A similar pattern is used in the EMM Plugin tests:

numba/numba/cuda/tests/cudadrv/test_emm_plugins.py

Lines 107 to 115 in 3497fb2

def setUp(self):

# Always start afresh with a new context and memory manager

cuda.close()

cuda.set_memory_manager(DeviceOnlyEMMPlugin)

def tearDown(self):

# Unset the memory manager for subsequent tests

cuda.close()

cuda.cudadrv.driver._memory_manager = None

(note this is not exactly the same because of the way the EMM Plugin interface works, but the idea is that the state for the test is set up in setUp() and then a suitable state for subsequent tests restored in tearDown()).

gmarkall · 2021-04-06T10:59:22Z

+        with override_config('LOW_OCCUPANCY_WARNINGS', 1):
+            with warnings.catch_warnings(record=True) as w:
+                warnings.simplefilter('always', NumbaPerformanceWarning)
+                ignore_internal_warnings()


What does ignore_internal_warnings() do here? I tried removing it and the test still seems to pass.

gmarkall · 2021-04-06T11:00:15Z

        copy2d(A, out=B)
        self.assertTrue(np.allclose(A, B))

+    def test_inefficient_guvectorize(self):


I think this could do with a comment explaining that the test creates an inefficient use of the GPU because the inputs are all mapped onto a single thread in a single block.

(Note: this is resolved, but I can't resolve the comment because I don't have write access to this repo)

gmarkall · 2021-04-06T11:04:00Z

+        smcount = ctx.device.MULTIPROCESSOR_COUNT
+
+        if (config.LOW_OCCUPANCY_WARNINGS) and (num_blocks < 2 * smcount):
+            msg = ("Number of blocks ({blocks}) that are  < 2 * SM ({sm}) "


I think will generate inefficient kernel code is a bit confusing - how about phrasing the warning as Grid size (<size>) less than 2 * SM count (2 * (<SM count>) = (<SM count * 2>) will likely result in GPU underutilization due to low occupancy?

… off

testhound · 2021-04-08T20:03:23Z

@gmarkall I believe all your review comments have been incorporated.

… in constructor

stuartarchibald

Thanks for the patch, this is great feedback to give users, it's good to see this implemented. I've a few minor comments in addition to @gmarkall's review, once resolved looks good to merge. Thanks again!

stuartarchibald · 2021-04-09T16:03:50Z


+        # Whether to warn about kernel launches where the grid size will
+        # underutilize the GPU due to low occupancy. On by default.
+        LOW_OCCUPANCY_WARNINGS = _readenv(


Suggested change

LOW_OCCUPANCY_WARNINGS = _readenv(

CUDA_LOW_OCCUPANCY_WARNINGS = _readenv(

All other CUDA related config vars start with CUDA_ or have CUDA in their name, suggest calling this something like the above.

stuartarchibald · 2021-04-09T16:04:21Z

+
+    def setUp(self):
+        self._low_occupancy_warnings = config.LOW_OCCUPANCY_WARNINGS
+        # Disable warnings about low gpu utilization for test suite


Suggested change

# Disable warnings about low gpu utilization for test suite

# Disable warnings about low gpu utilization in the test suite

stuartarchibald · 2021-04-09T16:06:22Z

+            grid_size = griddim[0] * griddim[1] * griddim[2]
+            if grid_size < 2 * smcount:
+                msg = ("Grid size ({grid}) < 2 * SM count ({sm}) "
+                       "will likely result in GPU underutilization due"


Suggested change

"will likely result in GPU underutilization due"

"will likely result in GPU under utilization due"

stuartarchibald · 2021-04-09T16:06:48Z

        )

+        # Whether to warn about kernel launches where the grid size will
+        # underutilize the GPU due to low occupancy. On by default.


Suggested change

# underutilize the GPU due to low occupancy. On by default.

# under utilize the GPU due to low occupancy. On by default.

stuartarchibald · 2021-04-09T16:08:36Z

+.. envvar:: NUMBA_CUDA_LOW_OCCUPANCY_WARNINGS
+
+   Enable warnings if the grid size is too small relative to the number of
+   steraming multiprocessors (SM). This option is on by default (1).


Suggested change

steraming multiprocessors (SM). This option is on by default (1).

streaming multiprocessors (SM). This option is on by default (default value is 1).

stuartarchibald · 2021-04-09T16:09:49Z

+   Enable warnings if the grid size is too small relative to the number of
+   steraming multiprocessors (SM). This option is on by default (1).
+
+   The heuristic checked is whether gridsize < 2 * (number of SMs). NOTE: The absence of


Suggested change

The heuristic checked is whether gridsize < 2 * (number of SMs). NOTE: The absence of

The heuristic checked is whether ``gridsize < 2 * (number of SMs)``. NOTE: The absence of

stuartarchibald · 2021-04-09T16:10:50Z

+   a warning does not imply a good gridsize relative to the number of SMs. Disabling
+   this warning will reduce the number of CUDA API calls, as the heuristic needs to
+   check the number of SMs for the device in the current context.


Does lead to a noticeable compilation speed gain?

This is for launch speed rather than compilation speed. In the past there have been performance issues due to excessive getting of contexts, e.g.: #5189 - so I think if we add a get of the context in the launch pipeline, it's good to have a way around it as well.

gmarkall · 2021-04-12T22:11:14Z

@gmarkall I believe all your review comments have been incorporated.

Thanks for looking at these. I think they are all addressed, except this one (from the initial review):

I think there should also be a test of this using a normal kernel launch - I can't see a really good place to put this new test, but either in numba/cuda/tests/cudapy/test_errors.py or in a new file test_warnings.py.

testhound · 2021-04-13T05:25:13Z

@gmarkall I believe all your review comments have been incorporated.

Thanks for looking at these. I think they are all addressed, except this one (from the initial review):

I think there should also be a test of this using a normal kernel launch - I can't see a really good place to put this new test, but either in numba/cuda/tests/cudapy/test_errors.py or in a new file test_warnings.py.

@gmarkall I addressed this with the function 'test_efficient_guvectorize' in the file test_gufunc.py.

gmarkall · 2021-04-13T22:48:10Z

Sorry, I forgot to respond earlier - I mean a test where you call a kernel like f[1,1](arg) and check the message is produced (and one that launches with better occupancy and checks that there is no warning). I do think test_efficient_guvectorize is a good addition too though.

gmarkall

Thanks for the updated tests! I have a few comments regarding the kernel launch test - rather than annotating lines on the diff, I'll try and summarise here:

The tests are quite complex compared to what's needed, resulting in the runtime being too long and the purpose of the tests being slightly obscured - the runtime on my machine is 40 seconds for test_warnings.py. Since we're only trying to test whether the warning is generated for a low occupancy configuration, we don't need the kernel to do anything or create any arrays to pass into it.
I think that even if the test runtime is reduced a lot, the test still doesn't work in the CUDA simulator, because it doesn't generate performance warnings (and we shouldn't try to make it do so).
The names of the tests are a bit vague about what is really tested - inefficient_kernel implies that the problem is inefficient code in the kernel, rather than the launch configuration.
A minor point: I try to not include CUDA in the names of CUDA test classes, because it's redundant in the CUDA test suite - the main issue I have with it is having to type more when running individual classes from the command line :-)

I think the above can be resolved with the following implementation of the test class:

@skip_on_cudasim('cudasim does not raise performance warnings')
class TestWarnings(CUDATestCase):
    def test_inefficient_launch_configuration(self):
        @cuda.jit
        def kernel():
            pass

        with override_config('CUDA_LOW_OCCUPANCY_WARNINGS', 1):
            with warnings.catch_warnings(record=True) as w:
                kernel[1, 1]()

        self.assertEqual(w[0].category, NumbaPerformanceWarning)
        self.assertIn('Grid size', str(w[0].message))
        self.assertIn('2 * SM count', str(w[0].message))

    def test_efficient_launch_configuration(self):
        @cuda.jit
        def kernel():
            pass

        with override_config('CUDA_LOW_OCCUPANCY_WARNINGS', 1):
            with warnings.catch_warnings(record=True) as w:
                kernel[256, 256]()

        self.assertEqual(len(w), 0)

There are also a couple of comments / questions on the diff to look at - one is related to the way setUp works, which turned out to be more complex than I first imagined it would be.

gmarkall · 2021-04-22T10:16:46Z

+    def __init__(self, *args, **kwargs):
+        super().__init__(*args, **kwargs)
+        self._low_occupancy_warnings = config.CUDA_LOW_OCCUPANCY_WARNINGS


I first thought this __init__() method should be entirely superfluous, because the setUp() method sets self._low_occupancy_warnings, and is called before every test. However, when I removed it and ran the test suite, many tests errored in tearDown() due to self._low_occupancy_warnings not being set. This is because some CUDATestCase subclasses also define setUp(), which replaces the setUp() defined here. I think the right way round this is for the setUpClass() class method to ensure that for those classes, both the CUDATestCase.setUp() method and their setUp() method are called. The following can be used to ensure this:

Suggested change

def __init__(self, *args, **kwargs):

super().__init__(*args, **kwargs)

self._low_occupancy_warnings = config.CUDA_LOW_OCCUPANCY_WARNINGS

@classmethod

def setUpClass(cls):

# If the test class defines a setUp method, we need to run both

# CUDATestCase.setUp (to disable low occupancy warnings) and its setUp

if cls is not CUDATestCase and cls.setUp is not CUDATestCase.setUp:

cls_setUp = cls.setUp

def setUp_both(self, *args, **kwargs):

# Call our setUp first so that low occupancy warnings are

# disabled before anything else happens

CUDATestCase.setUp(self)

return cls_setUp(self, *args, **kwargs)

cls.setUp = setUp_both

@gmarkall good catch on setUpClass. I don't think however, removing the __init __ method and it's implementation is correct. If you remove init then 'self._low_occupancy_warnings" will not be created as a instance variable.

It will be created during the setUp() method - it doesn't need to be "declared" in __init__().

gmarkall · 2021-04-22T10:41:14Z

+
+        with override_config('CUDA_LOW_OCCUPANCY_WARNINGS', 1):
+            with warnings.catch_warnings(record=True) as w:
+                warnings.simplefilter('always', NumbaPerformanceWarning)


What does this line do? (similar applies to all lines with warnings.simplefilter('always', NumbaPerformanceWarning) - I can't see that they have any effect)

gmarkall · 2021-04-22T11:02:54Z

        copy2d(A, out=B)
        self.assertTrue(np.allclose(A, B))

+    def test_inefficient_guvectorize(self):


(Note: this is resolved, but I can't resolve the comment because I don't have write access to this repo)

gmarkall · 2021-04-22T11:15:19Z

+        a = np.random.rand(1024 * 1024 * 32).astype('float32')
+        b = np.random.rand(1024 * 1024 * 32).astype('float32')


Looking at the performance of the kernel test led me to also look into how long the guvectorize tests take - I think the data sizes for this and the test below are too large for unit tests, and they take about 5 seconds to run - can we reduce the size of the data please - e.g.:

Suggested change

a = np.random.rand(1024 * 1024 * 32).astype('float32')

b = np.random.rand(1024 * 1024 * 32).astype('float32')

a = np.random.rand(1024 * 32).astype('float32')

b = np.random.rand(1024 * 32).astype('float32')

gmarkall · 2021-04-22T11:16:40Z

+        a = np.random.rand(1024 * 1024 * 128).astype('float32').\
+            reshape((1024 * 1024, 128))
+        b = np.random.rand(1024 * 1024 * 128).astype('float32').\
+            reshape((1024 * 1024, 128))


Also to reduce the data size here please:

Suggested change

a = np.random.rand(1024 * 1024 * 128).astype('float32').\

reshape((1024 * 1024, 128))

b = np.random.rand(1024 * 1024 * 128).astype('float32').\

reshape((1024 * 1024, 128))

a = np.random.rand(524288 * 2).astype('float32').reshape((524288, 2))

b = np.random.rand(524288 * 2).astype('float32').reshape((524288, 2))

The leading dimension of 524288 is needed to get a big enough grid size for a large GPU, so I reduced the trailing dimension down to 2 to keep the data size as small as possible.

testhound · 2021-04-22T21:51:24Z

@gmarkall should the tests in test_gufunc be skipped for the CUDA simulator (since no warnings are generated in the simulator), similar to the new ones in test_warning.py?

… class setUp

gmarkall

Thanks for the updates - looks good to me!

@stuartarchibald Do you want another buildfarm run, or are you happy with this as-is?

gmarkall · 2021-05-19T15:54:43Z

... Actually - if a buildfarm run can be squeezed in just to be on the safe side, that would be good - is that possible? (@stuartarchibald / @esc)

sklam

Tests are failing because of the following:

sklam · 2021-05-21T16:17:26Z

    Tests the DeviceRecord class with np.record host types
    """
    def setUp(self):
+        CUDATestCase.setUp()


Missing self in the current spelling. My suggestion is to use super().setUp() instead.

This will call TestCudaDeviceRecord.setUp(), which is what I was trying to avoid with the suggestion to use CUDATestCase.setUp() instead.

In that case, it needs to be CUDATestCase.setUp(self)

Example of errors:

====================================================================== ERROR: test_auto_device (numba.cuda.tests.cudadrv.test_cuda_devicerecord.TestCudaDeviceRecordWithRecord) ---------------------------------------------------------------------- Traceback (most recent call last): File "/Users/siu/dev/numba/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py", line 89, in setUp CUDATestCase.setUp() TypeError: setUp() missing 1 required positional argument: 'self' ====================================================================== ERROR: test_device_record_copy (numba.cuda.tests.cudadrv.test_cuda_devicerecord.TestCudaDeviceRecordWithRecord) ---------------------------------------------------------------------- Traceback (most recent call last): File "/Users/siu/dev/numba/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py", line 89, in setUp CUDATestCase.setUp() TypeError: setUp() missing 1 required positional argument: 'self' ====================================================================== ERROR: test_device_record_interface (numba.cuda.tests.cudadrv.test_cuda_devicerecord.TestCudaDeviceRecordWithRecord) ---------------------------------------------------------------------- Traceback (most recent call last): File "/Users/siu/dev/numba/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py", line 89, in setUp CUDATestCase.setUp() TypeError: setUp() missing 1 required positional argument: 'self' ====================================================================== ERROR: test_from_record_like (numba.cuda.tests.cudadrv.test_cuda_devicerecord.TestCudaDeviceRecordWithRecord) ---------------------------------------------------------------------- Traceback (most recent call last): File "/Users/siu/dev/numba/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py", line 89, in setUp CUDATestCase.setUp() TypeError: setUp() missing 1 required positional argument: 'self' ----------------------------------------------------------------------

Apologies for missing self in my earlier suggestion - I thought I tested my suggestions locally, but I must have done something not quite right!

gmarkall

LGTM!

stuartarchibald

Thanks for the patch and fixes @testhound.

Thanks for reviewing @gmarkall.

stuartarchibald · 2021-05-25T09:38:37Z

Buildfarm ID: numba_smoketest_cuda_yaml_68.

stuartarchibald · 2021-05-25T10:37:37Z

Buildfarm ID: numba_smoketest_cuda_yaml_68.

Passed

testhound added 4 commits March 29, 2021 15:52

Initial commit for guvectorize warning.

2ead91e

Add test cases and fix import issue

8fe4d4e

Merge branch 'master' into testhound/cuda_guvectorize

9d7336c

Add documentation for NUMBA_CUDA_LOW_OCCUPANCY_WARNINGS and use it

d8e74ab

testhound requested review from esc, sklam and stuartarchibald as code owners April 2, 2021 04:29

sklam added 3 - Ready for Review CUDA CUDA related issue/PR Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm labels Apr 2, 2021

gmarkall requested changes Apr 6, 2021

View reviewed changes

gmarkall added 4 - Waiting on author Waiting for author to respond to review and removed 3 - Ready for Review labels Apr 6, 2021

gmarkall added this to the Numba 0.54 RC milestone Apr 6, 2021

testhound added 3 commits April 8, 2021 12:30

Add new testcase, fix env var issue, reduce CUDA API calls if env var…

435fb91

… off

Merge branch 'master' into testhound/cuda_guvectorize

7412c6c

Remove outdated doc comment.

867167b

Fix build errors and declare _low_occupancy_warnings instance varible…

50d590e

… in constructor

stuartarchibald reviewed Apr 9, 2021

View reviewed changes

testhound added 2 commits April 12, 2021 22:00

Address upstream doc comments, and env var naming

e0fc84d

Merge branch 'master' into testhound/cuda_guvectorize

a5a0e4a

testhound added 2 commits April 21, 2021 23:25

Add new file to test kernel efficiency warnings

2a0544b

Merge branch 'master' into testhound/cuda_guvectorize

3724b8e

gmarkall requested changes Apr 22, 2021

View reviewed changes

gmarkall added 4 - Waiting on author Waiting for author to respond to review and removed 4 - Waiting on CI Review etc done, waiting for CI to finish 4 - Waiting on reviewer Waiting for reviewer to respond to author labels May 19, 2021

Remove unnecessary setUp call and add setUp call to CUDATestCase base…

3d1df9b

… class setUp

gmarkall previously approved these changes May 19, 2021

View reviewed changes

gmarkall added 4 - Waiting on CI Review etc done, waiting for CI to finish and removed 4 - Waiting on author Waiting for author to respond to review labels May 19, 2021

gmarkall added Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm and removed BuildFarm Passed For PRs that have been through the buildfarm and passed labels May 19, 2021

sklam requested changes May 21, 2021

View reviewed changes

sklam added the 4 - Waiting on author Waiting for author to respond to review label May 21, 2021

testhound added 2 commits May 21, 2021 16:23

Call CUDATestCase.setUp correctly

a86f553

Merge branch 'master' into testhound/cuda_guvectorize

d8e92be

testhound dismissed gmarkall’s stale review via d8e92be May 22, 2021 00:09

gmarkall approved these changes May 24, 2021

View reviewed changes

gmarkall added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 4 - Waiting on author Waiting for author to respond to review labels May 24, 2021

Merge branch 'master' into testhound/cuda_guvectorize

548b2f8

stuartarchibald approved these changes May 25, 2021

View reviewed changes

stuartarchibald removed the 4 - Waiting on reviewer Waiting for reviewer to respond to author label May 25, 2021

sklam approved these changes May 25, 2021

View reviewed changes

sklam merged commit 4bcea52 into numba:master May 25, 2021

testhound deleted the testhound/cuda_guvectorize branch June 1, 2021 23:13


		.. envvar:: NUMBA_CUDA_LOW_OCCUPANCY_WARNINGS

		Enable warnings if the block size is too small relative to the number of

		# whether to warn about kernels where the block size will
		# result in a inefficient mapping to GPUs. On by default.

	def setUp(self):
	# Always start afresh with a new context and memory manager
	cuda.close()
	cuda.set_memory_manager(DeviceOnlyEMMPlugin)

	def tearDown(self):
	# Unset the memory manager for subsequent tests
	cuda.close()
	cuda.cudadrv.driver._memory_manager = None

	LOW_OCCUPANCY_WARNINGS = _readenv(
	CUDA_LOW_OCCUPANCY_WARNINGS = _readenv(

	# Disable warnings about low gpu utilization for test suite
	# Disable warnings about low gpu utilization in the test suite

	"will likely result in GPU underutilization due"
	"will likely result in GPU under utilization due"

	# underutilize the GPU due to low occupancy. On by default.
	# under utilize the GPU due to low occupancy. On by default.

	steraming multiprocessors (SM). This option is on by default (1).
	streaming multiprocessors (SM). This option is on by default (default value is 1).

	The heuristic checked is whether gridsize < 2 * (number of SMs). NOTE: The absence of
	The heuristic checked is whether ``gridsize < 2 * (number of SMs)``. NOTE: The absence of

-    def __init__(self, *args, **kwargs):
-        super().__init__(*args, **kwargs)
-        self._low_occupancy_warnings = config.CUDA_LOW_OCCUPANCY_WARNINGS
+    @classmethod
+    def setUpClass(cls):
+        # If the test class defines a setUp method, we need to run both
+        # CUDATestCase.setUp (to disable low occupancy warnings) and its setUp
+        if cls is not CUDATestCase and cls.setUp is not CUDATestCase.setUp:
+            cls_setUp = cls.setUp
+            def setUp_both(self, *args, **kwargs):
+                # Call our setUp first so that low occupancy warnings are
+                # disabled before anything else happens
+                CUDATestCase.setUp(self)
+                return cls_setUp(self, *args, **kwargs)
+            cls.setUp = setUp_both

		a = np.random.rand(1024 * 1024 * 32).astype('float32')
		b = np.random.rand(1024 * 1024 * 32).astype('float32')

Conversation

testhound commented Apr 2, 2021

Uh oh!

sklam commented Apr 2, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

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

testhound commented Apr 8, 2021

Uh oh!

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

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gmarkall commented Apr 12, 2021

Uh oh!

testhound commented Apr 13, 2021

Uh oh!

gmarkall commented Apr 13, 2021

Uh oh!

gmarkall 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

sklam commented Apr 2, 2021 •

edited

Loading

sklam May 21, 2021 •

edited

Loading