[SYCL] Implement device libraries: C assert #922

asavonic · 2019-12-10T11:18:14Z

Device libraries provide a way to support functions from standard C
and C++ system headers in SYCL device code.

Runtime library requirements (functions that are normally defined in
libc, libm or libstdc++) are provided in a SPIR-V
library (libsycl-fallback.spv), which is linked at JIT (or AOT) time.

If a particular device supports a library OpenCL
extension (e.g. cl_intel_devicelib_assert), then SPIR-V implementation
is not linked and the device compiler has to provide definitions for
the corresponding functions.

NOTE:
The patch is a bit rough in program_manager.cpp after the recent refactoring changes there.
I'm going to refactor this bit soon.

Signed-off-by: Andrew Savonichev [email protected]

Device libraries provide a way to support functions from standard C and C++ system headers in SYCL device code. Runtime library requirements (functions that are normally defined in libc, libm or libstdc++) are provided in a SPIR-V library (libsycl-fallback.spv), which is linked at JIT (or AOT) time. If a particular device supports a library OpenCL extension (e.g. cl_intel_devicelib_assert), then SPIR-V implementation is not linked and the device compiler has to provide definitions for the corresponding functions. Signed-off-by: Andrew Savonichev <[email protected]>

sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst

sycl/include/CL/sycl/detail/context_impl.hpp

Signed-off-by: Andrew Savonichev <[email protected]>

Since build() may return a different program object, we need to properly release the old one (or not, if we return the same object). Signed-off-by: Andrew Savonichev <[email protected]>

Signed-off-by: Andrew Savonichev <[email protected]>

This should fix: fatal error: 'CL/__spirv/spirv_vars.hpp' file not found Signed-off-by: Andrew Savonichev <[email protected]>

sycl/source/detail/context_impl.cpp

sycl/source/detail/devicelib/glibc_wrapper.cpp

sycl/source/detail/devicelib/msvc_wrapper.cpp

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

sycl/include/CL/sycl/detail/program_manager/program_manager.hpp

sycl/source/detail/devicelib/CMakeLists.txt

sycl/source/detail/program_manager/program_manager.cpp

sycl/test/devicelib/assert-windows.cpp

Signed-off-by: Andrew Savonichev <[email protected]>

Clang calls it under the hood, but it has no explicit dependency for it in cmake. Signed-off-by: Andrew Savonichev <[email protected]>

Signed-off-by: Andrew Savonichev <[email protected]>

romanovvlad

LGTM

sycl/source/detail/devicelib/CMakeLists.txt

sycl/source/detail/devicelib/fallback-cassert.cpp

romanovvlad · 2019-12-16T12:26:03Z

sycl/source/detail/program_manager/program_manager.cpp

+  PI_CALL(piContextGetInfo)(Context, PI_CONTEXT_INFO_NUM_DEVICES,
+                            sizeof(NumDevices), &NumDevices, nullptr);
+  Devices.resize(NumDevices);
+  PI_CALL(piContextGetInfo)(Context, PI_CONTEXT_INFO_DEVICES,


Minor. Please handle possible error from piContextGetInfo

PI_CALL is supposed to throw an exception if the call fails, right? Not sure what else is missing.

sycl/source/detail/program_manager/program_manager.cpp

romanovvlad · 2019-12-16T14:07:17Z

sycl/source/detail/program_manager/program_manager.cpp

+    RT::PiResult Error = PI_CALL_NOCHECK(piProgramBuild)(
+        Program.get(), Devices.size(), Devices.data(), Opts, nullptr, nullptr);
+    if (Error != PI_SUCCESS)
+      compile_program_error(getProgramBuildLog(Program.get()));


OMG, we really need some static analysis tool here!

romanovvlad · 2019-12-16T14:10:42Z

sycl/source/detail/program_manager/program_manager.cpp

+      // an incomplete (but valid) LinkedProg, and throw.
+      throw compile_program_error(getProgramBuildLog(LinkedProg));
+    }
+    pi::checkPiResult(Error);


What happens here? Probably throw more user friendly exception than checkPiResult throws here?

Or at least leave a TODO

Current check should throw an exception and mention the error code. What should I mention in addition to that?

romanovvlad · 2019-12-16T14:12:28Z

sycl/source/detail/program_manager/program_manager.cpp

+    }
+    pi::checkPiResult(Error);
+  }
+  return Program;


Suggest checking that error == CL_SUCCESS and return so error handling is not in the if body.

It is going to look weird:

if (Error != PI_SUCCESS) { return Program; } if (LinkedProg) { // A non-trivial error occurred during linkage: get a build log, release // an incomplete (but valid) LinkedProg, and throw. throw compile_program_error(getProgramBuildLog(LinkedProg)); } pi::checkPiResult(Error);

A function that returns non-void ends with just pi::checkPiResult(Error);.
It is not obvious that the function is going to throw an exception.

sycl/test/devicelib/assert-windows.cpp

sycl/include/CL/sycl/detail/context_impl.hpp

sycl/source/detail/program_manager/program_manager.cpp

sergey-semenov

Grammar nitpicks are not blocking, of course.

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst

sycl/source/detail/program_manager/program_manager.cpp

sycl/test/devicelib/assert.cpp

sycl/source/detail/program_manager/program_manager.cpp

Signed-off-by: Andrew Savonichev <[email protected]>

sycl/source/detail/program_manager/program_manager.cpp

Signed-off-by: Andrew Savonichev <[email protected]>

sycl/source/detail/program_manager/program_manager.cpp

sycl/test/devicelib/assert.cpp

Signed-off-by: Andrew Savonichev <[email protected]>

AlexeySachkov

Overall, looks good to me, I don't have any major comments.

Anyway, I would like to @sergey-semenov to approve this patch too

keryell

The global idea of allowing more normal code in kernels on accelerators is great.
But some points are unclear to me.
Probably I miss some internal context.

sycl/doc/SYCLEnvironmentVariables.md

keryell · 2019-12-19T01:39:03Z

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

+       cgh.parallel_for<class SimpleVadd<T>>(numOfItems,
+       [=](cl::sycl::id<1> wiID) {
+           accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
+           assert(accessorC[wiID] > 0 && "Invalid value");


Is assert a good example? It relies on some kernel features not available in OpenCL... Are there real use cases?
What about this implementation:

#ifdef __SYCL_DEVICE_ONLY__ #define NDEBUG 1 #endif

? :-)

Is assert a good example? It relies on some kernel features not available in OpenCL...

Right, assert is an edge case and it is probably a bad example.
However, it is the only standard C and C++ function supported at this moment.

Are there real use cases?

For assert? It is widely used as a sanity check in host code, so I don't see a reason why it cannot be used for device code as well.

What about this implementation:

#ifdef __SYCL_DEVICE_ONLY__ #define NDEBUG 1 #endif

That should work if you define the macro before #include <assert.h>.
Beware any significant changes under #if NDEBUG:

struct shared { int i; #ifndef NDEBUG int debug; #endif };

I am unsure that assert() is really "supported", in the sense that the behaviour of what is proposed here is undefined.

The macro __SYCL_DEVICE_ONLY__ is set by the SYCL device compiler.
Having the device compiler driver to also set NDEBUG would be abusive but would somehow solve a little bit of the "problem".
But you could always #define assert() in your sycl.hpp even after including "cassert".
Anyway all this looks messy because, well, this is outside of the scope of SYCL and kernels...

Having the device compiler driver to also set NDEBUG would be abusive but would somehow solve a little bit of the "problem".
But you could always #define assert() in your sycl.hpp even after including "cassert".

There are two problems:

This solution depends on order of #include directives: if you include assert.h after sycl.hpp, the macro will be redefined again.

#ifdef _ASSERT_H # undef _ASSERT_H # undef assert [...] #endif /* assert.h */

Even if you manage to define the macro properly, there is still a problem of not having a proper abort function.

So we avoid this dance around __assert_fail and _wassert, but we still have to manage the devicelib assert extension.

Anyway all this looks messy because, well, this is outside of the scope of SYCL and kernels...

Agree.

keryell · 2019-12-19T01:40:26Z

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

+libraries, and allows to use them in SYCL device code.
+
+Function declarations are taken from the standard headers (e.g. from
+<assert.h> or <complex>), and the corresponding header has to be


<assert.h> or <complex>

But note that normally you should use in C++ instead...

keryell · 2019-12-19T01:41:04Z

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

+   clang++ -fsycl main.obj %SYCL_INSTALL%/lib/libsycl-msvc.o -o a.exe
+
+List of supported functions from C standard library:
+  - assert macro (from assert.h)


keryell · 2019-12-19T01:45:04Z

sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst

+and it indicates that a program does not execute as expected.
+The function should print a message containing the information
+provided in the arguments. In addition to that, the function is free
+to terminate the current kernel invocation.


And if not, what happens next?

It returns to the caller, and the kernel continues.

Even if it launch a nuclear missile that should have been avoided?

I would prefer the program not to compile at the first place!!!

I'd advise against using SYCL for anything that involves radioactive materials or weapons of mass destruction.

Even if it launch a nuclear missile that should have been avoided?

Wait, have you guarded this just by an assertion? Not the best idea, IMO

sycl/source/detail/devicelib/glibc_wrapper.cpp

sycl/source/detail/program_manager/program_manager.cpp

keryell · 2019-12-19T02:10:27Z

sycl/test/devicelib/assert.cpp

+// assert() call in device code guarantees nothing: on some devices it behaves
+// in the usual way and terminates the program. On other devices it can print an
+// error message and *continue* execution. Less capable devices can even ignore
+// an assert!


So the conclusion is...

So the conclusion is...

... the assert's behavior in SYCL device code is unspecified. However, you should get the expected behavior on devices that have capabilities to do that.

OK. But how to do this in a consistent and reasonable way?

You have to implement cl_intel_devicelib_cassert extension for all devices you care about to guarantee consistency between them.

keryell · 2019-12-19T02:12:09Z

sycl/test/devicelib/assert.cpp

+//
+//   Fallback mode (aka the best we can do by following the OpenCL spec):
+//     1. Assertion condition is printed to *stdout* by the OpenCL printf().
+//     2. Process (both host and device) is terminated by a SIGSEGV.


I am lost here. The fallback in SYCL is just the normal C++ code on the host so I have the feeling it should be the same for a "normal" assert() of the system.
What is the expectation of the end-user? That some functions are just changed by the Intel implementation when in kernels on the host? How a programmer can be sure that we can call the right function?

The fallback in SYCL is just the normal C++ code on the host so I have the feeling it should be the same for a "normal" assert() of the system.

The fallback in this context is a generic implementation of standard library functions provided in a SPIR-V form. This implementation is not aware about device-specific features, and can only use what is available in SPIR-V (i.e. no termination is possible for assert).

If a device compiler can do something better than a generic implementation, it can override some functions to achieve performance or functional improvements.

Anyway, the line 20 is misleading (2. Process is terminated). The SIGSEGV is not a side-effect from the fallback implementation: an unreachable instruction follows an __assert_fail call, so the process terminates after reaching it.

keryell · 2019-12-19T02:17:43Z

sycl/test/devicelib/assert.cpp

+//
+//   Native mode (same behavior as libc assert on CPU):
+//     1. Assertion condition is printed to *stderr*.
+//     2. Process (both host and device) is terminated by a SIGABRT.


So the fallback mode is not the native mode?

Or are you thinking to some OpenCL native kernels that could be used in SYCL interoperability mode or even by some OpenCL compiler for x86? https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clEnqueueNativeKernel.html

So the fallback mode is not the native mode?

Basically, we have not 2, but 3 modes:

Host device, where everything works already (it is a standard C++ code after all)

Fallback mode (on OpenCL/PI device), which is supposed to work with any device

Native mode (on OpenCL/PI device) provides the best possible implementation for a particular device.

OK, this seems different from what OpenCL defines as "native" then.

Signed-off-by: Andrew Savonichev <[email protected]>

sergey-semenov

LGTM overall, just one concern. sycl::program and sycl::kernel classes, generally speaking, cannot be used for kernels with assert() as of this patch. Are there plans to address that? This limitation should probably be mentioned somewhere.

asavonic · 2019-12-20T14:30:17Z

LGTM overall, just one concern. sycl::program and sycl::kernel classes, generally speaking, cannot be used for kernels with assert() as of this patch. Are there plans to address that? This limitation should probably be mentioned somewhere.

This is because they don't use functions from program manager, right?
Any ideas why?

sergey-semenov · 2019-12-20T14:41:58Z

This is because they don't use functions from program manager, right?
Any ideas why?

Right, current program manager API does not cover all possible use-cases (e. g. program::compile_with_program_type). It will have to be extended and then used in sycl::program instead of direct calls to PI.

bader · 2019-12-24T14:49:14Z

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

+For example, on Linux with GNU glibc:
+.. code:
+   clang++ -fsycl -c main.cpp -o main.o
+   clang++ -fsycl main.o $(SYCL_INSTALL)/lib/libsycl-glibc.o -o a.out


I think device library should be linked by default i.e. the same way OpenMP/CUDA compiler links device side library + driver option to disable linking device code.

Right. This requires changes in clang driver.

bader · 2019-12-24T14:51:48Z

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

+
+Implementation requires a special device library to be linked with a
+SYCL program. The library should match the C or C++ standard library
+used to compile the program:


The library should match the C or C++ standard library used to compile the program:

Can we have a single device library?

No, see the line 77 in this document.
Basically, different libraries may use the same function name (e.g. __assert_fail), but with different arguments.

I don't think it's the case for existing C/C++ standard library implementations.
Assuming that we want avoid potential issues with conflicting names, why do you propose "a special device library" instead of "a special header" as it's done for CUDA/OpenMP?

I don't think it's the case for existing C/C++ standard library implementations.

These little inconsistencies definitely exist.
For example, in musl C library, __assert_fail has an line argument of type int, and Glibc has this argument as unsigned int.

Even if we consider only one library (e.g. Glibc), I'm not sure that ABI between headers and the library will always remain stable, so future/past versions may bring similar problems.

Assuming that we want avoid potential issues with conflicting names, why do you propose "a special device library" instead of "a special header" as it's done for CUDA/OpenMP?

Can you describe this "special header" approach using the assert function as an example?

Please, take a look here: https://github.com/intel/llvm/blob/sycl/clang/lib/Headers/__clang_cuda_runtime_wrapper.h#L335.
I suggest reading comments in https://github.com/intel/llvm/blob/sycl/clang/lib/Headers/__clang_cuda_* and take a look at https://github.com/intel/llvm/tree/sycl/clang/lib/Headers/cuda_wrappers.

It is hard to tell how this approach is going to co-exist with the rest of the standard library, and scale beyond a handful of functions. The code you refer to doesn't seem to support anything other than glibc and libc++/libstdc++.

In general, I think the most difficult part here is (1) to establish an interface with a device compiler (DevicelibExtensions.rst), and (2) wire standard library functions through that interface, instead of using library specific routines.

(2) can probably be done in a header file, but I don't see a huge benefit in comparison with the library approach. If you (or anyone reading this) see a benefit and know to to implement this with portability in mind - feel free to send a patch.

I think this at least should be discussed with the clang community.
It should be safe to develop solution aligned with the existing approaches in the community code w/o constantly checking with the community, but in this case we went in slightly different direction.

bader · 2019-12-24T15:03:14Z

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

+
+.. code:
+   clang++ -fsycl -c main.cpp -o main.o
+   clang++ -fsycl main.o $(SYCL_INSTALL)/lib/libsycl-glibc.o -o a.out


Two points:

Do we really need a separate wrapper for each standard library implementation? I don't think I understand the problem with a boilerplate code in *every* device compiler. I think linker is able to import only necessary functions, so scattering "boilerplate" code into separate wrapper libraries seems unnecessary.

If we use single device library, it can be linked by the driver implicitly, but even if we stick to existing proposal, driver is able to infer right library from the host triple.

(1) should be explained on line 77. Let me know if I should elaborate.
(2) can be done as a separate change in driver.

(1) I understand that we should have definitions for __assert_fail, _wassert and __assert_func to support Glibc, MSVC and newlib, but can we keep them all in the same library? What is the strategy for the symbols shared between different implementation (if any): e.g. Glibc and newlib?

(1) I understand that we should have definitions for __assert_fail, _wassert and __assert_func to support Glibc, MSVC and newlib, but can we keep them all in the same library?

This doesn't work for functions that have the same name, but different arguments (e.g. __assert_fail in glibc vs musl).

What is the strategy for the symbols shared between different implementation (if any): e.g. Glibc and newlib?

There are no shared symbols. If there is a function with identical name and signature, it will be duplicated in every wrapper library.

In any case, I anticipate the wrapper libraries to go away soon. @s-kanaev is working on a header-based approach to provide the wrappers.

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

asavonic requested a review from romanovvlad December 10, 2019 11:23

asavonic assigned AlexeySachkov Dec 10, 2019

agozillon reviewed Dec 10, 2019

View reviewed changes

sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst Outdated Show resolved Hide resolved

agozillon reviewed Dec 10, 2019

View reviewed changes

sycl/include/CL/sycl/detail/context_impl.hpp Outdated Show resolved Hide resolved

agozillon reviewed Dec 10, 2019

View reviewed changes

sycl/include/CL/sycl/detail/context_impl.hpp Outdated Show resolved Hide resolved

Andrew Savonichev added 12 commits December 11, 2019 16:52

Do not use "user" options for device library compilation

6a8b3e7

Signed-off-by: Andrew Savonichev <[email protected]>

Pass unique_ptr with a program to ProgramManager::build()

bc335d7

Since build() may return a different program object, we need to properly release the old one (or not, if we return the same object). Signed-off-by: Andrew Savonichev <[email protected]>

Refactoring

b03b3e2

Signed-off-by: Andrew Savonichev <[email protected]>

Clang format the whole patch

08a37e3

Signed-off-by: Andrew Savonichev <[email protected]>

Code style change

7b57b6b

Signed-off-by: Andrew Savonichev <[email protected]>

Release library programs

462af28

Signed-off-by: Andrew Savonichev <[email protected]>

Cleanup invalid library programs if an exception is thrown

52e077c

Signed-off-by: Andrew Savonichev <[email protected]>

Tabs vs spaces

75c05b5

Signed-off-by: Andrew Savonichev <[email protected]>

Reword documentation in context_impl.hpp

269d6dc

Signed-off-by: Andrew Savonichev <[email protected]>

Fix error handling of clLinkProgram

3d80fe0

Signed-off-by: Andrew Savonichev <[email protected]>

Code style changes

7497a23

Signed-off-by: Andrew Savonichev <[email protected]>

Include path to SYCL headers

82e0e84

This should fix: fatal error: 'CL/__spirv/spirv_vars.hpp' file not found Signed-off-by: Andrew Savonichev <[email protected]>

s-kanaev reviewed Dec 12, 2019

View reviewed changes

sycl/source/detail/context_impl.cpp Show resolved Hide resolved

s-kanaev reviewed Dec 12, 2019

View reviewed changes

sycl/source/detail/devicelib/glibc_wrapper.cpp Show resolved Hide resolved

sycl/source/detail/devicelib/msvc_wrapper.cpp Show resolved Hide resolved

AlexeySachkov reviewed Dec 12, 2019

View reviewed changes

Andrew Savonichev added 3 commits December 13, 2019 17:14

Re-format PI_CALL

265aaf6

Signed-off-by: Andrew Savonichev <[email protected]>

Misc code review comments

4ce3a57

Signed-off-by: Andrew Savonichev <[email protected]>

Misc code review comments

278206f

Signed-off-by: Andrew Savonichev <[email protected]>

AlexeySachkov requested a review from sergey-semenov December 13, 2019 14:35

Andrew Savonichev added 2 commits December 13, 2019 20:34

Add a dependency for llvm-spirv

a5d65db

Clang calls it under the hood, but it has no explicit dependency for it in cmake. Signed-off-by: Andrew Savonichev <[email protected]>

Disable the test for assert on Windows

6b67601

Signed-off-by: Andrew Savonichev <[email protected]>

romanovvlad reviewed Dec 16, 2019

View reviewed changes

AlexeySachkov reviewed Dec 16, 2019

View reviewed changes

sycl/source/detail/program_manager/program_manager.cpp Outdated Show resolved Hide resolved

sycl/source/detail/program_manager/program_manager.cpp Show resolved Hide resolved

sycl/source/detail/program_manager/program_manager.cpp Outdated Show resolved Hide resolved

sergey-semenov reviewed Dec 16, 2019

View reviewed changes

sycl/source/detail/program_manager/program_manager.cpp Outdated Show resolved Hide resolved

Andrew Savonichev added 2 commits December 17, 2019 18:01

Reword documentation

d585921

Signed-off-by: Andrew Savonichev <[email protected]>

Use 'if WIN32' instead of 'if MSVC' in CMake

e882460

Signed-off-by: Andrew Savonichev <[email protected]>

dm-vodopyanov reviewed Dec 17, 2019

View reviewed changes

sycl/source/detail/program_manager/program_manager.cpp Outdated Show resolved Hide resolved

Andrew Savonichev added 9 commits December 17, 2019 21:18

Code review comments

f1ae75f

Signed-off-by: Andrew Savonichev <[email protected]>

Get devices from context instead of going through OpenCL API

a1aeda7

Signed-off-by: Andrew Savonichev <[email protected]>

Refactor getDeviceLibPrograms

f3a4e07

Signed-off-by: Andrew Savonichev <[email protected]>

Use an enum instead of std::string for extension id

27bcf97

Signed-off-by: Andrew Savonichev <[email protected]>

Fix typo

45b9be8

Signed-off-by: Andrew Savonichev <[email protected]>

Merge branch 'intel/sycl' into private/asavonic/assert

73fdcda

Use detail::get_device_info to get device extensions

9f1aecc

Signed-off-by: Andrew Savonichev <[email protected]>

Fix typo

8e826ef

Signed-off-by: Andrew Savonichev <[email protected]>

Use uint64_t instead of size_t for __devicelib_assert_fail

cf819f1

Signed-off-by: Andrew Savonichev <[email protected]>

AlexeySachkov reviewed Dec 18, 2019

View reviewed changes

Andrew Savonichev added 2 commits December 18, 2019 18:04

Code review comments

8b41e4a

Signed-off-by: Andrew Savonichev <[email protected]>

Workaround bug in IGC

423a054

Signed-off-by: Andrew Savonichev <[email protected]>

AlexeySachkov previously approved these changes Dec 18, 2019

View reviewed changes

keryell reviewed Dec 19, 2019

View reviewed changes

Andrew Savonichev added 2 commits December 19, 2019 19:06

Re-order local and global id in the assert format string

42d52b8

Signed-off-by: Andrew Savonichev <[email protected]>

Code review comments

2f61342

Signed-off-by: Andrew Savonichev <[email protected]>

asavonic dismissed AlexeySachkov’s stale review via 2f61342 December 19, 2019 16:11

sergey-semenov reviewed Dec 20, 2019

View reviewed changes

romanovvlad approved these changes Dec 23, 2019

View reviewed changes

romanovvlad merged commit 0039ee0 into intel:sycl Dec 24, 2019

bader reviewed Dec 24, 2019

View reviewed changes

asavonic mentioned this pull request Dec 25, 2019

[clang][Bundler] Add a CMake dependency for llvm-objcopy #970

Merged

[SYCL] Implement device libraries: C assert #922

[SYCL] Implement device libraries: C assert #922

Uh oh!

Conversation

asavonic commented Dec 10, 2019

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

romanovvlad left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

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!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

sergey-semenov left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

AlexeySachkov left a comment

Choose a reason for hiding this comment

Uh oh!

keryell left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

asavonic Dec 20, 2019 •

edited

Loading