Add HiveHash support for nested types #9

ustcfy · 2024-10-11T05:37:34Z

From this link, we can see the nested HiveHash computation in the Spark CPU version.

Here is a simple example of calculating the hash value for a List<List<int>> type.

Assuming column_device_view col is a List<List<int>> type:

cudf::column_device_view curr_col = col.slice(row_index, 1);
Now curr_col may look like the following:

Index	List<List<int>>
0	{{1, 0}, null, {2, null}}

You can run the following code in the Spark shell to calculate the hash value of curr_col.

import org.apache.spark.sql.catalyst.expressions._
import org.apache.spark.sql.types._
import org.apache.spark.sql.catalyst.util._

val inType = BoundReference(0, ArrayType(ArrayType(IntegerType)), true)
val hh = HiveHash(Seq(inType))

def toAData(any:Any): ArrayData = if(any==null) null else ArrayData.toArrayData(any)
 
val data = Seq(
    List(List(1, 0), null, List(2, null))
  )

data.map { row =>
 val rr = if (row == null) null else row.map(toAData(_))
 val r = toAData(rr)
 hh.eval(new GenericInternalRow(Array(r.asInstanceOf[Any])))
}

Signed-off-by: ustcfy <[email protected]>

firestarman · 2024-10-11T06:26:01Z

src/main/cpp/src/hive_hash.cu


-    template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
+    template <typename T>


We'd better keep the two original methods here to separate the non-nested types from nested types. It can avoid some useless "if" condition check and "accumulate" execution. Which will impact the perf when there are more than millions of rows.

firestarman · 2024-10-11T06:28:37Z

src/main/cpp/src/hive_hash.cu

    __device__ hive_hash_value_t operator()(cudf::column_device_view const& col,
                                            cudf::size_type row_index) const noexcept
    {
-      CUDF_UNREACHABLE("Nested type is not supported");
+      cudf::column_device_view curr_col = col.slice(row_index, 1);
+      stack_element stack[8] = {stack_element(curr_col), stack_element(curr_col), stack_element(curr_col), stack_element(curr_col),


Why is the size "8" ?

i can not understand how stack way works. Better to add some comment here.

Why is the size "8" ?

It's just a temporary solution, and I will improve it later.

Move the 8 into the template.

firestarman · 2024-10-11T06:32:53Z

src/test/java/com/nvidia/spark/rapids/jni/HashTest.java

+          true, false, null, false, true, null);
+         ColumnView structs = ColumnView.makeStructView(strings, integers, doubles, floats, bools);
+         ColumnVector result = Hash.hiveHash(new ColumnView[]{structs});
+         ColumnVector expected = Hash.hiveHash(new ColumnVector[]{strings, integers, doubles, floats, bools})) {


This can not test the correctness of HiveHash computation. It can only prove that nested struct has nothing to do with the hash result. If the hash algorithm produces different output than Spark HiveHash, the test can still pass.

res-life · 2024-10-11T06:48:41Z

src/test/java/com/nvidia/spark/rapids/jni/HashTest.java

+         ColumnView structs = ColumnView.makeStructView(strings, integers, doubles, floats, bools);
+         ColumnVector result = Hash.hiveHash(new ColumnView[]{structs});
+         ColumnVector expected = Hash.hiveHash(new ColumnVector[]{strings, integers, doubles, floats, bools})) {
+      assertColumnsAreEqual(expected, result);


Test cases are using the same interface to get both expected and actual results.
Please leverage the results in testHiveHashStrings, testHiveHashFloats...
Test case will like:

struct_data = {1.0, 2.3, 5} actual = Hash.hiveHash(...) expected = 37 * (37 * hash(1.0) + hash(2.3)) + hash(5) assert_equals(expected, actual)

It's better that we can use Spark hive hash interface to get the expected result.
Try to introduce the corresponding Java dependency with test scope.

<dependency> <groupId>xx</groupId> <artifactId>xx</artifactId> <scope>test</scope> </dependency>

I would not suggest do this in JNI repo. We can hardcode the expected results just as what other tests do.
But we can test them against Spark hive hash in Rapids tests.

thirtiseven

I tried to read the stack logic, it looks like something correct, but could we add some more comments?

~~Also if there is not much effort in plugin side, I suggest we can test there with integration test, which has a larger coverage and more convenient to write.~~ Ok seems it is not so convenient for this case

thirtiseven · 2024-10-11T06:04:22Z

src/main/cpp/src/hive_hash.cu

    __device__ hive_hash_value_t operator()(cudf::column_device_view const& col,
                                            cudf::size_type row_index) const noexcept
    {
-      CUDF_UNREACHABLE("Nested type is not supported");
+      cudf::column_device_view curr_col = col.slice(row_index, 1);
+      stack_element stack[8] = {stack_element(curr_col), stack_element(curr_col), stack_element(curr_col), stack_element(curr_col),


Question: Does this mean that we only support nested types whose depth is less than 8?

Yes, should fallback in the Plugin when encounter a nest type is large than 8 depths.

res-life · 2024-10-11T08:51:22Z

src/test/java/com/nvidia/spark/rapids/jni/HashTest.java

@@ -510,4 +510,164 @@ void testHiveHashMixed() {
      assertColumnsAreEqual(expected, result);
    }
  }
+
+  @Test
+  void testHiveHashStruct() {


Would you please test all the types, especially DICTIONARY32, TIMESTAMP_MILLISECONDS, DECIMAL32 types.

switch (dtype.id()) { case type_id::INT8: return f.template operator()<typename IdTypeMap<type_id::INT8>::type>( std::forward<Ts>(args)...); case type_id::INT16: return f.template operator()<typename IdTypeMap<type_id::INT16>::type>( std::forward<Ts>(args)...); case type_id::INT32: return f.template operator()<typename IdTypeMap<type_id::INT32>::type>( std::forward<Ts>(args)...); case type_id::INT64: return f.template operator()<typename IdTypeMap<type_id::INT64>::type>( std::forward<Ts>(args)...); case type_id::UINT8: return f.template operator()<typename IdTypeMap<type_id::UINT8>::type>( std::forward<Ts>(args)...); case type_id::UINT16: return f.template operator()<typename IdTypeMap<type_id::UINT16>::type>( std::forward<Ts>(args)...); case type_id::UINT32: return f.template operator()<typename IdTypeMap<type_id::UINT32>::type>( std::forward<Ts>(args)...); case type_id::UINT64: return f.template operator()<typename IdTypeMap<type_id::UINT64>::type>( std::forward<Ts>(args)...); case type_id::FLOAT32: return f.template operator()<typename IdTypeMap<type_id::FLOAT32>::type>( std::forward<Ts>(args)...); case type_id::FLOAT64: return f.template operator()<typename IdTypeMap<type_id::FLOAT64>::type>( std::forward<Ts>(args)...); case type_id::BOOL8: return f.template operator()<typename IdTypeMap<type_id::BOOL8>::type>( std::forward<Ts>(args)...); case type_id::TIMESTAMP_DAYS: return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_DAYS>::type>( std::forward<Ts>(args)...); case type_id::TIMESTAMP_SECONDS: return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_SECONDS>::type>( std::forward<Ts>(args)...); case type_id::TIMESTAMP_MILLISECONDS: return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_MILLISECONDS>::type>( std::forward<Ts>(args)...); case type_id::TIMESTAMP_MICROSECONDS: return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_MICROSECONDS>::type>( std::forward<Ts>(args)...); case type_id::TIMESTAMP_NANOSECONDS: return f.template operator()<typename IdTypeMap<type_id::TIMESTAMP_NANOSECONDS>::type>( std::forward<Ts>(args)...); case type_id::DURATION_DAYS: return f.template operator()<typename IdTypeMap<type_id::DURATION_DAYS>::type>( std::forward<Ts>(args)...); case type_id::DURATION_SECONDS: return f.template operator()<typename IdTypeMap<type_id::DURATION_SECONDS>::type>( std::forward<Ts>(args)...); case type_id::DURATION_MILLISECONDS: return f.template operator()<typename IdTypeMap<type_id::DURATION_MILLISECONDS>::type>( std::forward<Ts>(args)...); case type_id::DURATION_MICROSECONDS: return f.template operator()<typename IdTypeMap<type_id::DURATION_MICROSECONDS>::type>( std::forward<Ts>(args)...); case type_id::DURATION_NANOSECONDS: return f.template operator()<typename IdTypeMap<type_id::DURATION_NANOSECONDS>::type>( std::forward<Ts>(args)...); case type_id::DICTIONARY32: return f.template operator()<typename IdTypeMap<type_id::DICTIONARY32>::type>( std::forward<Ts>(args)...); case type_id::STRING: return f.template operator()<typename IdTypeMap<type_id::STRING>::type>( std::forward<Ts>(args)...); case type_id::LIST: return f.template operator()<typename IdTypeMap<type_id::LIST>::type>( std::forward<Ts>(args)...); case type_id::DECIMAL32: return f.template operator()<typename IdTypeMap<type_id::DECIMAL32>::type>( std::forward<Ts>(args)...); case type_id::DECIMAL64: return f.template operator()<typename IdTypeMap<type_id::DECIMAL64>::type>( std::forward<Ts>(args)...); case type_id::DECIMAL128: return f.template operator()<typename IdTypeMap<type_id::DECIMAL128>::type>( std::forward<Ts>(args)...); case type_id::STRUCT: return f.template operator()<typename IdTypeMap<type_id::STRUCT>::type>( std::forward<Ts>(args)...);

@res-life

Do you mean, for example, to add fields like TIMESTAMP_MILLISECONDS and DECIMAL32 inside struct? However, it seems that the Hivehash for these two types has not been implemented yet.

As for DICTIONARY, I will consider adding support if the cases for struct and list are resolved.

Do you mean, for example, to add fields like TIMESTAMP_MILLISECONDS and DECIMAL32 inside struct? However, it seems that the Hivehash for these two types has not been implemented yet.

If Hivehash does not support these two types, we should try to find a way to implement it, or we should fallback in the Rapids plugin. I assume Spark supports these two types.

As for DICTIONARY, I will consider adding support if the cases for struct and list are resolved.
Please help test.

DICTIONARY for basic types like int, long

DICTIONARY for nested types, maybe cuDF can not generate this kind of data.

More primitive types support can be a follow-up. This PR focuses on nested types support.

res-life · 2024-10-12T03:28:38Z

src/main/cpp/src/hive_hash.cu

    __device__ hive_hash_value_t operator()(cudf::column_device_view const& col,
                                            cudf::size_type row_index) const noexcept
    {
-      CUDF_UNREACHABLE("Nested type is not supported");
+      cudf::column_device_view curr_col = col.slice(row_index, 1);
+      stack_element stack[8] = {stack_element(curr_col), stack_element(curr_col), stack_element(curr_col), stack_element(curr_col),


No need to initilize 8 column_device_view.
Consider:

constexpr int len_of_8_stack_element = 8 * sizeof(stack_element); // column_device_view default constructor is deleted, can not allocate column_device_view array directly // use byte array to wrapper stack_element list unit8 stack_wrapper[len_of_8_stack_element]; stack_element *stack = reinterpret_cast<stack_element*>(stack_wrapper);

Please add some pseudocode to explain the calculation logic:
e.g.:

// The stack is simulating the recursive function invokes: hive_hash(float) { return (int)float; } hive_hash(list) { int hash = 0 for(child : list) { // Depth first Traversal hash += hive_hash(child) } return hash }

It's better to add a simple example.

Signed-off-by: ustcfy <[email protected]>

…vehash-nested-support

Signed-off-by: ustcfy <[email protected]>

firestarman · 2024-10-14T02:10:49Z

src/main/cpp/src/hive_hash.cu

+      cudf::column_device_view curr_col = col.slice(row_index, 1);
+      // column_device_view default constructor is deleted, can not allocate column_device_view array directly
+      // use byte array to wrapper StackElement list
+      constexpr int len_of_8_StackElement = 8 * sizeof(StackElement);


why is 8 used here ?

src/main/cpp/src/hive_hash.cu

    template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
    __device__ hive_hash_value_t operator()(cudf::column_device_view const& col,
                                            cudf::size_type row_index) const noexcept
    {
-      CUDF_UNREACHABLE("Nested type is not supported");
+      hive_hash_value_t ret = HIVE_INIT_HASH;


firestarman · 2024-10-14T02:21:57Z

src/main/cpp/src/hive_hash.cu

+      uint8_t stack_wrapper[len_of_8_StackElement];
+      StackElementPtr stack = reinterpret_cast<StackElementPtr>(stack_wrapper);


I am confused on creating a stack of StackElement ?
Does the below not work in C++ ?

StackElement statck[8];

I am confused on creating a stack of StackElement ? Does the below not work in C++ ?

StackElement statck[8];

Because if I define an array StackElement stack[8]; without initialization, each element of the array will be constructed using the default constructor of StackElement. However, the default constructor for the member cudf::column_device_view in StackElement is deleted.

OK, But you can restore the default constructor. Any reason it has to be deleted?
uint8_t stack_wrapper[len_of_8_StackElement] will also initialize 8 empty StackElement, the same as what StackElement statck[8] will do.

Any reason it has to be deleted?

Do not know why cuDF marks default constructor as deleted.

column_device_view() = delete;

uint8_t stack_wrapper[len_of_8_StackElement] will also initialize 8 empty StackElement

For a array like A_Class[8], it will first allocate 8 * sizeof(A_Class) memory, then call default constructor to initilize 8 instences. uint8_t stack_wrapper[len_of_8_StackElement] will skip the call of default constructor to initilize.

Here is a workaround to skip the call of default constructor.

uint8_t stack_wrapper[len_of_8_StackElement]; StackElementPtr stack = reinterpret_cast<StackElementPtr>(stack_wrapper);

Signed-off-by: ustcfy <[email protected]>

firestarman · 2024-10-14T07:17:59Z

src/main/cpp/src/hive_hash.cu

@@ -157,7 +157,7 @@ hive_hash_value_t __device__ inline hive_hash_function<cudf::timestamp_us>::oper
 * @tparam hash_function Hash functor to use for hashing elements. Must be hive_hash_function.
 * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
 */
-template <template <typename> class hash_function, typename Nullate>
+template <template <typename> class hash_function, typename Nullate, int MAX_NESTED_LEN = 8>


Personally it is not a good idea to hardcode the nested length. It will crash when there are more than 8 nested layers.
If 8 is the upper limitation, we need to check the length and throw an exception (simliar as what check_hash_compatibility does)before running into this row hasher, and notice this limitation in the doc of this function/class.
Another feasible way is we can calculate this length from the schema (you can get the schema from the input table view) and allocate a GPU stack (e.g. rmm::device_vector, or rmm::device_buffer) and pass the stack into the hive_device_row_hasher.

Another feasible way is we can calculate this length from the schema (you can get the schema from the input table view) and allocate a GPU stack

get-json-object is also using a constant depth. I remember that allocate a GPU stack will impact performance.
I recommand that check the max length in CPU cpp code before call into this kernel and throw an exception if length exceeds 8. In plugin Java code, add fallback logic.

The GPU stack (actually a GPU memory) is allocated with a fixed size calculated from the schema, just like allocating an output column (a GPU memory inside). So it should have no big impact on perf, I think.

And personally this case is different from yours. This max length can be inferred from the schema without accessing any actual row. But get-json-object requires rows access to get this info.

res-life · 2024-10-14T07:49:06Z

src/main/cpp/src/hive_hash.cu

+            stack_size--;
+          } else {
+            // Push child column to stack
+            stack[stack_size++] = StackElement(cudf::detail::structs_column_device_view(curr_col).get_sliced_child(element->child_idx), element->factor_exp + curr_col.num_child_columns() - 1 - element->child_idx);


~~Add throw exception if stack_size > 8~~

Can we throw an exception in a CUDA kernel ?

See an example as below, a mcro is used to handle this kind of case. And CUDF_UNREACHABLE

template <typename T, CUDF_ENABLE_IF(not cudf::column_device_view::has_element_accessor<T>())> __device__ hash_value_type operator()(cudf::column_device_view const&, cudf::size_type, Nullate const, hash_value_type const) const noexcept { CUDF_UNREACHABLE("Unsupported type for xxhash64"); }

Throw an exception in CPU code before call into this kernel.

res-life · 2024-10-14T07:49:14Z

src/main/cpp/src/hive_hash.cu

+          if (element->child_idx == curr_col.size()) {
+            stack_size--;
+          } else {
+            stack[stack_size++] = StackElement(curr_col.slice(element->child_idx, 1), element->factor_exp + curr_col.size() - element->child_idx - 1);


Throw an exception in CPU code before call into this kernel.

Signed-off-by: ustcfy <[email protected]>

res-life · 2024-10-15T05:48:40Z

src/main/cpp/src/hive_hash.cu

@@ -157,7 +157,7 @@ hive_hash_value_t __device__ inline hive_hash_function<cudf::timestamp_us>::oper
 * @tparam hash_function Hash functor to use for hashing elements. Must be hive_hash_function.
 * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
 */
-template <template <typename> class hash_function, typename Nullate>
+template <template <typename> class hash_function, typename Nullate, int MAX_NESTED_LEN>


Nit:
Rename MAX_NESTED_LEN to MAX_NESTED_DEPTH

res-life · 2024-10-15T05:49:10Z

src/main/cpp/src/hive_hash.cu

@@ -239,6 +443,10 @@ std::unique_ptr<cudf::column> hive_hash(cudf::table_view const& input,
  // Return early if there's nothing to hash
  if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }

+  //Nested depth cannot exceed 8
+  int const max_nested_len = 8;


constexpr is better than int const

res-life · 2024-10-15T05:52:56Z

src/main/cpp/src/hive_hash.cu

+     * The hash value for the column is computed as:
+     * hash(S1) = hash(S2) * HIVE_HASH_FACTOR + hash(d)
+     *          = (hash(i) * HIVE_HASH_FACTOR + hash(f)) * HIVE_HASH_FACTOR + hash(d)
+     *          = hash(i) * HIVE_HASH_FACTOR^2 + hash(f) * HIVE_HASH_FACTOR + hash(d)


This transformation is not equivalent when considering overflow.

res-life · 2024-10-15T05:55:38Z

src/main/cpp/src/hive_hash.cu

+      stack[stack_size++] = StackElement(curr_col, 1);
+
+      while (stack_size > 0) {
+        StackElementPtr element = &stack[stack_size - 1];


Nit:
StackElement const& element = stack[stack_size - 1];

Nit: StackElement const& element = stack[stack_size - 1];

But I need to modify the child_idx member of the element.

res-life · 2024-10-15T06:01:27Z

src/main/cpp/src/hive_hash.cu

+        check_nested_depth_impl(*child, max_nested_len - 1);
+      }
+    } else {      // Primitive type
+      CUDF_EXPECTS(max_nested_len > 0, "The nested depth of the input table exceeds the maximum supported depth.");


Nit:
Add max_nested_len into error msg to tell report more details.

Signed-off-by: ustcfy <[email protected]>

src/main/cpp/src/hive_hash.cu

Co-authored-by: Liangcai Li <[email protected]>

src/main/cpp/src/hive_hash.cu

Co-authored-by: Liangcai Li <[email protected]>

firestarman · 2024-10-17T07:18:07Z

src/test/java/com/nvidia/spark/rapids/jni/HashTest.java

+  }
+
+  @Test
+  void testHiveHashNestedType() {


Better to have a test to verify the case when the max nested depth is larger than 8.
Also missing tests for the type of list of struct.

… depth limit Signed-off-by: ustcfy <[email protected]>

firestarman

Looks Good to Me

Signed-off-by: ustcfy <[email protected]>

Add HiveHash support for nested types

bf6e5e4

Signed-off-by: ustcfy <[email protected]>

firestarman reviewed Oct 11, 2024

View reviewed changes

res-life reviewed Oct 11, 2024

View reviewed changes

thirtiseven reviewed Oct 11, 2024

View reviewed changes

res-life reviewed Oct 11, 2024

View reviewed changes

res-life reviewed Oct 12, 2024

View reviewed changes

ustcfy added 3 commits October 12, 2024 16:28

Add HiveHash support for nested types

00ae762

Signed-off-by: ustcfy <[email protected]>

Merge remote-tracking branch 'ustcfy/hivehash-nested-support' into hi…

663bbd4

…vehash-nested-support

Resolve merge conflicts

35017e8

Signed-off-by: ustcfy <[email protected]>

ustcfy requested review from res-life, firestarman and thirtiseven October 12, 2024 09:30

firestarman reviewed Oct 14, 2024

View reviewed changes

Move hardcoded value 8 to template parameter in hive_device_row_hasher

1c987bf

Signed-off-by: ustcfy <[email protected]>

firestarman reviewed Oct 14, 2024

View reviewed changes

res-life reviewed Oct 14, 2024

View reviewed changes

Added more comments and nested level detection

9a18fae

Signed-off-by: ustcfy <[email protected]>

res-life reviewed Oct 15, 2024

View reviewed changes

Refactor code logic and hardcode expected results for tests.

32d6beb

Signed-off-by: ustcfy <[email protected]>

firestarman reviewed Oct 16, 2024

View reviewed changes

src/main/cpp/src/hive_hash.cu Outdated Show resolved Hide resolved

src/main/cpp/src/hive_hash.cu Outdated Show resolved Hide resolved

src/main/cpp/src/hive_hash.cu Outdated Show resolved Hide resolved

src/main/cpp/src/hive_hash.cu Outdated Show resolved Hide resolved

ustcfy and others added 4 commits October 16, 2024 17:02

Update src/main/cpp/src/hive_hash.cu

29a4204

Co-authored-by: Liangcai Li <[email protected]>

Update src/main/cpp/src/hive_hash.cu

d667003

Co-authored-by: Liangcai Li <[email protected]>

Update src/main/cpp/src/hive_hash.cu

e704199

Co-authored-by: Liangcai Li <[email protected]>

Update src/main/cpp/src/hive_hash.cu

9e98ce7

Co-authored-by: Liangcai Li <[email protected]>

firestarman reviewed Oct 16, 2024

View reviewed changes

src/main/cpp/src/hive_hash.cu Outdated Show resolved Hide resolved

Update src/main/cpp/src/hive_hash.cu

4813e4d

Co-authored-by: Liangcai Li <[email protected]>

firestarman reviewed Oct 17, 2024

View reviewed changes

Enhance stack_element encapsulation and add test for exceeding nested…

2824989

… depth limit Signed-off-by: ustcfy <[email protected]>

firestarman approved these changes Oct 21, 2024

View reviewed changes

ustcfy added 2 commits October 22, 2024 10:14

Format C++ code

499826f

Signed-off-by: ustcfy <[email protected]>

Format

c3d1b9a

Signed-off-by: ustcfy <[email protected]>

ustcfy marked this pull request as ready for review October 22, 2024 05:37

ustcfy changed the base branch from 0902 to 1010 October 22, 2024 06:46

Merge branch '1010' into hivehash-nested-support

56ff0e3

wjxiz1992 merged commit e3d5769 into wjxiz1992:1010 Oct 23, 2024
1 check passed

ustcfy deleted the hivehash-nested-support branch October 24, 2024 05:59

ustcfy mentioned this pull request Oct 24, 2024

Add HiveHash support for nested types NVIDIA/spark-rapids-jni#2534

Merged


		template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
		template <typename T>

		uint8_t stack_wrapper[len_of_8_StackElement];
		StackElementPtr stack = reinterpret_cast<StackElementPtr>(stack_wrapper);

Add HiveHash support for nested types #9

Add HiveHash support for nested types #9

Conversation

ustcfy commented Oct 11, 2024 • edited Loading

firestarman Oct 11, 2024 • edited Loading

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

thirtiseven left a comment • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

res-life Oct 11, 2024 • edited Loading

Choose a reason for hiding this comment

ustcfy Oct 11, 2024 • edited Loading

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

res-life Oct 12, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

This comment was marked as resolved.

Choose a reason for hiding this comment

Choose a reason for hiding this comment

firestarman Oct 14, 2024 • edited Loading

Choose a reason for hiding this comment

res-life Oct 14, 2024 • edited Loading

Choose a reason for hiding this comment

firestarman Oct 14, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

firestarman Oct 14, 2024 • edited Loading

Choose a reason for hiding this comment

res-life Oct 14, 2024 • edited Loading

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

res-life Oct 14, 2024 • edited Loading

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

firestarman Oct 17, 2024 • edited Loading

Choose a reason for hiding this comment

firestarman left a comment

Choose a reason for hiding this comment

ustcfy commented Oct 11, 2024 •

edited

Loading

firestarman Oct 11, 2024 •

edited

Loading

thirtiseven left a comment •

edited

Loading

res-life Oct 11, 2024 •

edited

Loading

ustcfy Oct 11, 2024 •

edited

Loading

res-life Oct 12, 2024 •

edited

Loading

firestarman Oct 14, 2024 •

edited

Loading

res-life Oct 14, 2024 •

edited

Loading

firestarman Oct 14, 2024 •

edited

Loading

firestarman Oct 14, 2024 •

edited

Loading

res-life Oct 14, 2024 •

edited

Loading

res-life Oct 14, 2024 •

edited

Loading

firestarman Oct 17, 2024 •

edited

Loading