Optimize K-Truss #4742

jnke2016 · 2024-11-01T04:23:31Z

This PR introduces several optimization to speed up K-Truss. In fact, our K-Truss implementation computes the intersection of all edges regardless they are weak or not which can be very expensive if only few edges need to be invalidated. By running nbr_intersection on the weak edges, this considerably improves the runtime.

…2_optimize-ktruss

ChuckHastings · 2024-11-21T21:22:02Z

cpp/src/community/k_truss_impl.cuh

+    auto has_edge = thrust::binary_search(
+          thrust::seq, edgelist_first, weak_edgelist_first, edge_p_r);
+
+    if (!has_edge) { // FIXME: Do binary search instead


Is this FIXME still relevant? It looks like you're doing binary search

ChuckHastings · 2024-11-21T21:24:39Z

cpp/src/community/k_truss_impl.cuh

+
+
+template <typename vertex_t, typename edge_t>
+struct extract_edges { // FIXME:  ******************************Remove this functor. For testing purposes only*******************


Should these be left in here?

ChuckHastings · 2024-11-21T21:27:49Z

cpp/src/community/k_truss_impl.cuh

    std::optional<rmm::device_uvector<weight_t>> edgelist_wgts{std::nullopt};
-
+
+    //#if 0


Delete the commented out lines in this section

Right I still had a lot of debug prints which are now removed

…2_optimize-ktruss

seunghwak · 2024-11-22T17:49:20Z

@jnke2016 Is this PR ready for review?

seunghwak · 2024-11-22T18:42:10Z

cpp/src/community/k_truss_impl.cuh

+
+{
+  // FIXME: Use global comm for debugging purposes
+  // then replace it by minor comm once the accuracy is verified


Haven't you verified this? Using global_comm vs minor_comm has a pretty significant scalability impact in large scale (even if you're sending to/receiving from the same set of ranks). All-to-all on a subcommunicator performs better than P2P on a subset of ranks using a global communicator.

seunghwak · 2024-11-22T18:54:17Z

cpp/src/community/k_truss_impl.cuh

+  std::vector<size_t> h_tx_counts(d_tx_counts.size());
+
+  raft::update_host(
+    h_tx_counts.data(), d_tx_counts.data(), d_tx_counts.size(), handle.get_stream());


handle.sync_stream() is necessary before using h_tx_counts.

No guarantee that h_tx_counts will hold valid values before handle.sync_stream().

Don't forget to address this.

seunghwak · 2024-11-22T18:54:59Z

cpp/src/community/k_truss_impl.cuh

+  std::tie(dsts, std::ignore) =
+    shuffle_values(handle.get_comms(), edgelist_dsts.begin(), h_tx_counts, handle.get_stream());
+
+  // rmm::device_uvector<bool> edge_exists(0, handle.get_stream());


seunghwak · 2024-11-22T19:09:25Z

cpp/src/community/k_truss_impl.cuh

+    shuffle_values(handle.get_comms(), edge_exists.begin(), rx_counts, handle.get_stream());
+
+  // The 'edge_exists' array is ordered based on 'cp_edgelist_srcs' where the edges where group,
+  // hoever it needs to match 'edgelist_srcs', hence re-order 'edge_exists' accordingly.


Is this true? I guess edge_exists is ordered based on edgelist_srcs not cp_edgelist_srcs.

Yes other way around. It is cp_edgelist_srcs that is supposed to be groupby_and_count. The documentation is correct but made a mistake in the group_by_and_count call. I don't want to mess up original edges in the triangles when grouping that's why I make a copy.

seunghwak · 2024-11-22T19:11:27Z

cpp/src/community/k_truss_impl.cuh

+void order_edge_based_on_dodg(raft::handle_t const& handle,
+                              graph_view_t<vertex_t, edge_t, false, multi_gpu>& graph_view,
+                              raft::device_span<vertex_t> edgelist_srcs,
+                              raft::device_span<vertex_t> edgelist_dsts)


What are you doing in this function?

Essentially this function is ordering the edges obtained from nbr_intersection (On the symmetric graph) based on the DODG edges

Once we identify weak edges, the next step is to retrieve the triangles incident to those weak edges and this is done with nbr_intersection (on the symmetric graph). Once we have the three endpoints of the triangle, the next step is to find the direction that matches the edges in the DODG and this is what that function does

…2_optimize-ktruss

seunghwak · 2024-11-27T00:06:43Z

cpp/src/community/k_truss_impl.cuh

+  std::vector<size_t> h_tx_counts(d_tx_counts.size());
+
+  raft::update_host(
+    h_tx_counts.data(), d_tx_counts.data(), d_tx_counts.size(), handle.get_stream());


Don't forget to address this.

seunghwak · 2024-11-27T03:30:38Z

cpp/src/community/k_truss_impl.cuh

+  std::optional<rmm::device_uvector<vertex_t>> cp_edgelist_dsts{std::nullopt};
+
+  // FIXME: Minor comm is not working for all cases so I believe some edges a beyong
+  // the partitioning range


a beyong here sounds like a typo.

And have you confirmed that some edge sources/destinations are outside the expected range? Or maybe the code to use minor_comm has a bug? Whichever is the case, we need to really dig into this and make sure everything is working as expected. This optimization is for large scale K-Truss benchmarking. Any testing/debugging/tuning gets way more difficult in larger scales. We need to verify everything we can in small scales before going to a larger scale. If we have doubts in our code even in this scale, things will get way worse in larger scales.

seunghwak · 2024-11-27T03:35:50Z

cpp/src/community/k_truss_impl.cuh

+
+    std::vector<size_t> h_tx_counts(d_tx_counts.size());
+
+    handle.sync_stream();


This should be after raft::update_host.

seunghwak · 2024-11-27T03:37:05Z

cpp/src/community/k_truss_impl.cuh

+      h_tx_counts.data(), d_tx_counts.data(), d_tx_counts.size(), handle.get_stream());
+
+    std::tie(srcs, rx_counts) = shuffle_values(
+      handle.get_comms(), cp_edgelist_srcs->begin(), h_tx_counts, handle.get_stream());


You have already created an alias for handle.get_comms() in line 66. Why use handle.get_comms() instead of just comm?

seunghwak · 2024-11-27T03:37:14Z

cpp/src/community/k_truss_impl.cuh

+      handle.get_comms(), cp_edgelist_srcs->begin(), h_tx_counts, handle.get_stream());
+
+    std::tie(dsts, std::ignore) = shuffle_values(
+      handle.get_comms(), cp_edgelist_dsts->begin(), h_tx_counts, handle.get_stream());


seunghwak · 2024-11-27T04:34:34Z

cpp/src/community/k_truss_impl.cuh

+                              raft::device_span<vertex_t> edgelist_srcs,
+                              raft::device_span<vertex_t> edgelist_dsts)
+
+{


So, here, what we are actually doing is this.

For (src, dst) pairs in edgelist_srcs & edgelist_dsts, return (src, dst) if (src, dst) exists in graph_view, and return (dst, src) if not.

I think this code is overly complex.

If single-GPU, this code looks fine.

In multi-GPU, you first need to find unique (src, dst) pairs, shuffle, call has_edge, shuffle_back. Now you have (src, dst, exist) triplets. Sort them for binary search. Then, you can iterate over the input pairs and flip or not based on whether the edge exists or not.

In the current code, you are finding unique edges after shuffling. But better do this at the beginning to reduce computation/communication. And for temporary variables, try to minimize their scopes. Some variables' scopes are too large and this makes code harder to understand.

seunghwak · 2024-11-27T04:35:38Z

cpp/src/community/k_truss_impl.cuh

+        auto src = thrust::get<0>(edgelist_first[idx]);
+        auto dst = thrust::get<1>(edgelist_first[idx]);
+
+        auto itr_pair = thrust::find(  // FIXME: replace by lower bound


This is a simple fix. Don't defer simple fixes for later updates.

…2_optimize-ktruss

seunghwak

I am reviewing the entire K-Truss code; more reviews coming.

seunghwak · 2024-12-04T19:41:31Z

cpp/src/community/k_truss_impl.cuh

@@ -214,121 +409,501 @@ k_truss(raft::handle_t const& handle,

  // 3. Keep only the edges from a low-degree vertex to a high-degree vertex.



auto [srcs, dsts, wgts] = k_core(handle, cur_graph_view, edge_weight_view, k - 1, std::make_optional(k_core_degree_type_t::OUT), std::make_optional(core_number_span));

I guess this code won't work if cur_graph_view != graph_view; edge_weight_view is invalid if cur_graph_view == modified_graph_view.

seunghwak · 2024-12-04T19:47:34Z

cpp/src/community/k_truss_impl.cuh

@@ -121,6 +314,8 @@ k_truss(raft::handle_t const& handle,
    edge_weight{std::nullopt};
  std::optional<rmm::device_uvector<weight_t>> wgts{std::nullopt};

+  cugraph::edge_bucket_t<vertex_t, void, true, multi_gpu, true> edgelist_dodg(handle);


Better minimize the scope of edgelist_dodg. This is used only for the transform_e in line 460. No reason to define this variable way ahead.

And I think maintaining both cur_graph_view and modified_graph_view is redundant.

Won't
auto cur_graph_view = modified_graph ? modified_graph.view() : graph_view be enough?

Right. Now I only maintain cur_graph_view and unmasked_cur_graph_view

Better minimize the scope of edgelist_dodg

After updating my script, this variable was no longer need and removed

seunghwak · 2024-12-04T19:55:17Z

cpp/src/community/k_truss_impl.cuh

-    }
-    renumber_map = std::move(tmp_renumber_map);
-  }
+  edgelist_dodg.clear();

  // 4. Compute triangle count using nbr_intersection and unroll weak edges

  {


What's the purpose of this curly bracket? This closes at the end of this function definition and doesn't help in reducing the scope of any variable.

seunghwak · 2024-12-04T19:57:07Z

cpp/src/community/k_truss_impl.cuh

@@ -214,121 +409,501 @@ k_truss(raft::handle_t const& handle,

  // 3. Keep only the edges from a low-degree vertex to a high-degree vertex.

-  {
-    auto cur_graph_view = modified_graph_view ? *modified_graph_view : graph_view;


We may place this DODG extraction part inside { ... }. We just need to define dodg_mask outside the bracket and I assume all the other variables aren't used besides extracting a DODG graph.

This was also addressed

seunghwak · 2024-12-04T20:11:16Z

cpp/src/community/k_truss_impl.cuh

@@ -121,6 +314,8 @@ k_truss(raft::handle_t const& handle,
    edge_weight{std::nullopt};


Should we maintain edge_weight from here? We can treat the graph as an unweighted graph till we finish finding K-Truss edge sources & destinations. Once we have K-Truss edge sources & destinations, we can extract edge (src, dst, weight) triplets from the input graph.

Right. I treated the graph as unweighted and extracted the edges with weights(if exist) at the end.

ChuckHastings · 2024-12-04T21:01:10Z

Moving this to 25.02... hopefully merging very early.

…2_optimize-ktruss

ChuckHastings · 2025-01-22T23:29:01Z

cpp/src/community/k_truss_impl.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2024, NVIDIA CORPORATION.


Looks like your copyright reverted... should include 2025.

ChuckHastings · 2025-01-22T23:34:37Z

cpp/src/community/k_truss_impl.cuh

@@ -41,8 +42,68 @@
 #include <thrust/transform.h>
 #include <thrust/tuple.h>

+#include <chrono>


Is this used?

Rather not have dangling using statements like this in a header. If we're actually using std::chrono, I think it's better to prefix the function calls/variable references (e.g. std::chrono::duration) rather than a using which loses the context of where duration comes from.

If this is just from some debugging/benchmarking that's not in use, best not to include it in the final code.

If this is just from some debugging/benchmarking that's not in use, best not to include it in the final code.

Yes those were only for benchmarking purposes and I removed them

…2_optimize-ktruss

seunghwak · 2025-01-28T04:21:29Z

cpp/src/community/k_truss_impl.cuh

+template <typename vertex_t, typename edge_t>
+struct extract_weak_edges {
+  edge_t k{};
+  __device__ thrust::optional<thrust::tuple<vertex_t, vertex_t>> operator()(


thrust::optional will be replaced with cuda::std::optional (see #4891).

While the above PR will replace thrust::optional in the old code, it won't replace thrust::optional in the new code. We should use cuda::std::optional in the new code.

Right. I updated my PR. As a matter of fact, I couldn't even build my branch once I merged the latest changes

seunghwak · 2025-01-28T04:22:14Z

cpp/src/community/k_truss_impl.cuh

+  __device__ thrust::optional<thrust::tuple<vertex_t, vertex_t>> operator()(
+    vertex_t src, vertex_t dst, thrust::nullopt_t, thrust::nullopt_t, edge_t count) const
+  {
+    return ((count < k - 2) && (count > 0))


Why should we check count > 0? A comment explaining this will be helpful.

Why should we check count > 0? A comment explaining this will be helpful.

count can only be >= 0 if we unrolled the edge count properly (without overcompensating). We don't need to process edges with count == 0 that's why I check count > 0. I just replaced that check by ((count < k - 2) && (count != 0)) which makes more sense.

Once I am done processing all the edges in the k-truss, I simply mask all the edges with count ==0 as you can see below and then extract the remaining edges with is the final result.

cugraph::transform_e( handle, cur_graph_view, cugraph::edge_src_dummy_property_t{}.view(), cugraph::edge_dst_dummy_property_t{}.view(), edge_triangle_counts.view(), [] __device__(auto src, auto dst, cuda::std::nullopt_t, cuda::std::nullopt_t, auto count) { return count == 0 ? false : true; }, dodg_mask.mutable_view(), false);

seunghwak · 2025-01-28T04:24:08Z

cpp/src/community/k_truss_impl.cuh

+};
+
+template <typename edge_t>
+struct is_k_minus_1_or_greater_t {


Isn't this a misnomer?

If core number is k - 1, this function will return false.

Should we rename to "is_k_or_greater_t"?

seunghwak · 2025-01-28T05:25:20Z

cpp/src/community/k_truss_impl.cuh

+    auto endpoints = thrust::make_tuple(weak_srcs[chunk_start + idx],  // p
+                                        weak_dsts[chunk_start + idx],  // q
+                                        intersection_indices[i]        // r
+    );
+
+    // Re-order the endpoints such that p < q < r in order to identify duplicate triangles
+    // which will cause overcompensation. comparing the vertex IDs is cheaper than comparing the
+    // degrees (d(p) < d(q) < d(r)) which will be done once in the latter stage to retrieve the
+    // direction of the edges once the triplet dependency is broken.
+    if (thrust::get<0>(endpoints) > thrust::get<2>(endpoints)) {  // (a > c)
+      endpoints = thrust::make_tuple(thrust::get<2>(endpoints),
+                                     thrust::get<1>(endpoints),
+                                     thrust::get<0>(endpoints));  // swap(a, c)
+    }
+
+    if (thrust::get<0>(endpoints) > thrust::get<1>(endpoints)) {  // (a > b)
+      endpoints = thrust::make_tuple(thrust::get<1>(endpoints),
+                                     thrust::get<0>(endpoints),
+                                     thrust::get<2>(endpoints));  // swap(a, b);
+    }
+
+    if (thrust::get<1>(endpoints) > thrust::get<2>(endpoints)) {  // (b > c)
+      endpoints = thrust::make_tuple(thrust::get<0>(endpoints),
+                                     thrust::get<2>(endpoints),
+                                     thrust::get<1>(endpoints));  // swap(b, c);
+    }


I think you can simplify this.

#include <cuda/std/utility> ... auto p = weak_srcs[chunk_start + idx]; auto q = weak_dsts[chunk_start + idx]; auto r = intersectoin_indices[i]; // Re-order the endpoints such that p < q < r in order to identify duplicate triangles // which will cause overcompensation. comparing the vertex IDs is cheaper than comparing the // degrees (d(p) < d(q) < d(r)) which will be done once in the latter stage to retrieve the // direction of the edges once the triplet dependency is broken. if (p > q) cuda::std::swap(p, q); if (p > r) cuda::std::swap(p, r); if (q > r) cuda::std::swap(q, r); return std::make_tuple(p, q, r);

seunghwak · 2025-01-28T06:55:23Z

cpp/src/community/k_truss_impl.cuh

-    }
+    // 2.1 Exclude self-loops
+
+    if (graph_view.count_self_loops(handle) > edge_t{0}) {


Better not mix graph_view & cur_graph_view (even though they are same at this point).

if (cur_graph_view.count_self_loops(handle) > edge_t{0}) {

could be less confusing.

seunghwak · 2025-01-28T06:59:35Z

cpp/src/community/k_truss_impl.cuh

+      cugraph::edge_property_t<decltype(cur_graph_view), bool> in_k_minus_1_core_edge_mask(
+        handle, cur_graph_view);
+      cugraph::fill_edge_property(
+        handle, unmasked_cur_graph_view, in_k_minus_1_core_edge_mask.mutable_view(), false);


It is desirable to minimize the scope of a variable. Better move this right before the transform_e call to update in_k_minus_1_core_edge_mask.

seunghwak · 2025-01-28T07:03:00Z

cpp/src/community/k_truss_impl.cuh

+      dodg_mask.mutable_view(),
+      false);


dodg_mask.mutable_view());

No need to explicitly set do_expensive_check to false here.

Right. But since we have the expensive check exposed, I can just pass it to every function that support it.

seunghwak · 2025-01-28T07:07:54Z

cpp/src/community/k_truss_impl.cuh

+      edge_src_dummy_property_t{}.view(),
+      edge_dst_dummy_property_t{}.view(),
+      edge_mask.view(),
+      [] __device__(auto src, auto dst, auto src_out_degree, auto dst_out_degree, auto e_mask) {


auto src_out_degree, auto dst_out_degree,

=>

auto, auto,

src_out_degree & dst_out_degree are misnomers... They are dummy values.

Right, I am not passing anything for the src and dat property.

seunghwak · 2025-01-28T07:27:59Z

cpp/src/community/k_truss_impl.cuh

+        return e_mask ? true : false;
+      },
+      weak_edges_mask.mutable_view(),
+      false);


No need to explicitly set do_expensive_check to false here.

Won't weak_edge_mask = std::move(edge_mask) work here? (Or may just rename edge_mask to something like undirected_mask? So, we just maintain undirected_mask & dodg_mask.)

weak_edge_mask = std::move(edge_mask)

Good catch. No need to repeat the same work in edge_mask with transform_e

seunghwak · 2025-01-28T07:33:16Z

cpp/src/community/k_truss_impl.cuh

+        return count == 0 ? false : true;
+      },
+      dodg_mask.mutable_view(),
+      false);


No need to explicitly set do_expensive_check = false.

seunghwak · 2025-01-28T07:37:37Z

cpp/src/community/k_truss_impl.cuh

+               minor_comm_size}] __device__(auto val) {
+            return key_func(thrust::get<0>(val), thrust::get<1>(val));
+          },
+


No need for an empty line here.

seunghwak

LGTM

ChuckHastings · 2025-01-30T20:00:01Z

/merge

jnke2016 added 3 commits October 31, 2024 19:30

optimize ktruss

e935857

benchmark k-truss

e61a580

add benchmark print

0c83b54

github-actions bot added the cuGraph label Nov 1, 2024

add weights support

2e0ef1a

jnke2016 self-assigned this Nov 15, 2024

jnke2016 added this to the 24.12 milestone Nov 15, 2024

jnke2016 marked this pull request as ready for review November 18, 2024 15:32

jnke2016 requested a review from a team as a code owner November 18, 2024 15:32

jnke2016 added 2 commits November 21, 2024 05:30

Merge remote-tracking branch 'upstream/branch-24.12' into branch-24.1…

89597c5

…2_optimize-ktruss

update SG implementation of K-Truss and add MG

1259d41

ChuckHastings reviewed Nov 21, 2024

View reviewed changes

jnke2016 added 5 commits November 21, 2024 22:01

add function to reorder the edges based on the DODG

7f48a91

Merge remote-tracking branch 'upstream/branch-24.12' into branch-24.1…

13a70cb

…2_optimize-ktruss

remove debug print statement

9ad6bfd

fix style

b515cd9

Merge remote-tracking branch 'upstream/branch-24.12' into branch-24.1…

895cb0e

…2_optimize-ktruss

ChuckHastings approved these changes Nov 22, 2024

View reviewed changes

seunghwak reviewed Nov 22, 2024

View reviewed changes

jnke2016 added 3 commits November 26, 2024 16:36

add sync call and fix typos

3dfae34

fix style

8fa9cd9

Merge remote-tracking branch 'upstream/branch-24.12' into branch-24.1…

7b8999d

…2_optimize-ktruss

seunghwak reviewed Nov 27, 2024

View reviewed changes

jnke2016 added 2 commits December 3, 2024 17:24

reduce before shuffling

106afb3

Merge remote-tracking branch 'upstream/branch-24.12' into branch-24.1…

9cabbd2

…2_optimize-ktruss

seunghwak reviewed Dec 4, 2024

View reviewed changes

ChuckHastings modified the milestones: 24.12, 25.02 Dec 4, 2024

jnke2016 added 2 commits January 22, 2025 10:18

Merge remote-tracking branch 'upstream/branch-25.02' into branch-24.1…

5a1be1d

…2_optimize-ktruss

removebenchmark print

b98ddb0

ChuckHastings changed the base branch from branch-24.12 to branch-25.02 January 22, 2025 20:27

ChuckHastings reviewed Jan 22, 2025

View reviewed changes

create dodg without extracting the edges

cd4a104

github-actions bot removed CMake python ci benchmarks conda labels Jan 24, 2025

jnke2016 added 3 commits January 27, 2025 15:20

fix style

ee1d1f0

Merge remote-tracking branch 'upstream/branch-25.02' into branch-24.1…

8ee8a07

…2_optimize-ktruss

remove benchmark import

c7c96ae

seunghwak reviewed Jan 28, 2025

View reviewed changes

ChuckHastings removed request for a team and gforsyth January 28, 2025 16:35

jnke2016 added 7 commits January 29, 2025 07:57

replace thrust::optional by cuda::std::optional

394bb62

update check

9ac3965

simplify script

53a0ba1

reduce variable scope

105d08a

rename variables and simplify code

a91c6d5

fix style

3fa3468

update branch

1dfbad2

seunghwak approved these changes Jan 29, 2025

View reviewed changes

ChuckHastings added improvement Improvement / enhancement to an existing function non-breaking Non-breaking change labels Jan 30, 2025

rapids-bot bot merged commit 4baef58 into rapidsai:branch-25.02 Jan 30, 2025
81 checks passed



		template <typename vertex_t, typename edge_t>
		struct extract_edges { // FIXME: ****************************Remove this functor. For testing purposes only*****************

		std::optional<rmm::device_uvector<weight_t>> edgelist_wgts{std::nullopt};


		//#if 0


		std::vector<size_t> h_tx_counts(d_tx_counts.size());

		handle.sync_stream();

		@@ -214,121 +409,501 @@ k_truss(raft::handle_t const& handle,

		// 3. Keep only the edges from a low-degree vertex to a high-degree vertex.

		@@ -121,6 +314,8 @@ k_truss(raft::handle_t const& handle,
		edge_weight{std::nullopt};

Optimize K-Truss #4742

Optimize K-Truss #4742

Conversation

jnke2016 commented Nov 1, 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

seunghwak commented Nov 22, 2024

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

jnke2016 Nov 22, 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

seunghwak left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

ChuckHastings commented Dec 4, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jnke2016 Jan 29, 2025 • 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

seunghwak left a comment

Choose a reason for hiding this comment

ChuckHastings commented Jan 30, 2025

jnke2016 commented Nov 1, 2024 •

edited

Loading

jnke2016 Nov 22, 2024 •

edited

Loading

jnke2016 Jan 29, 2025 •

edited

Loading