Add local to global index mapping #1707

MarcelKoch · 2024-10-23T15:12:31Z

This PR adds a local-to-global mapping to the index_map class. It is necessary for the follow-up PR #1639. This also introduces a device-side view of a partition, and adds another member to the partition which stores the range-ids segmented by their part-id.

(I know that the branch name is wrong...)

yhmtsai

not finish review yet.
why is it required for PGM?
after merging row and column, do they still need to map back to the global index?
you need to reindex all of them because the size will be the smaller than the original one after coarsening.

yhmtsai · 2024-12-18T08:09:16Z

core/distributed/device_partition.hpp

+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GINKGO_PARTITION_HPP


nit:

Suggested change

#ifndef GINKGO_PARTITION_HPP

#ifndef GKO_CORE_DISTRIBUTED_DEVICE_PARTITION_HPP_

I think it is probably time now to move to #pragma once ?

core/distributed/device_partition.hpp

core/distributed/index_map_kernels.hpp

reference/distributed/partition_kernels.cpp

test/distributed/index_map_kernels.cpp

yhmtsai · 2024-12-18T12:14:55Z

core/distributed/partition.cpp

+    exec->run(partition::make_build_ranges_by_part(
+        part_ids_.get_const_data(), get_num_ranges(), get_num_parts(),
+        range_ids, num_ranges_per_part));
+    ranges_by_part_ = segmented_array<size_type>::create_from_sizes(
+        std::move(range_ids), num_ranges_per_part);


currently, this rely on the atomic add to count the size.
then segmented_array will convert size to offset.
Maybe you can implement a similar kernel with convert_idxs_to_ptrs + additional index map from i to range_id[i] because range_part[range_id[i]] is sorted. then segmented_array can create_from_offsets directly

It can be future improvement not in this pr

Yes, for now I would just use atomics. I don't think this is a performance critical part, so it should be fine for a while.

reference/distributed/partition_helpers.hpp

MarcelKoch · 2024-12-18T12:52:03Z

after merging row and column, do they still need to map back to the global index?
you need to reindex all of them because the size will be the smaller than the original one after coarsening.

@yhmtsai I don't understand your point, could you rephrase this?

yhmtsai · 2024-12-18T13:00:20Z

rank 0: holds row 0, 1, 4
rank 1: holds row 2, 3
if we mege 0, 1 and 2, 3
rank 0: 0, 4
rank 1: 2
we need to compute the size of each gap from different node and then reindex the global index
rank 0: 0, 2
rank 1: 1

MarcelKoch · 2024-12-18T13:02:37Z

@yhmtsai I think this discussion makes more sense in #1639. I don't see much relevance to this PR.

yhmtsai · 2024-12-18T13:04:11Z

yes, but if PGM does not need it, do you still need this function for something else?

MarcelKoch · 2024-12-18T13:08:25Z

I think I need at least parts of it to keep the old distributed matrix constructor around. @fritzgoebel also seems to need it for his domain decomposition matrix.

Anyway, this still sounds like a discussion for #1639.

pratikvn

Some comments, but mostly looks good!

pratikvn · 2025-01-22T10:54:05Z

core/distributed/device_partition.hpp

+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GINKGO_PARTITION_HPP


I think it is probably time now to move to #pragma once ?

core/distributed/device_partition.hpp

pratikvn · 2025-01-22T11:47:19Z

core/distributed/device_partition.hpp

+    return {num_parts,
+            partition->get_num_empty_parts(),
+            partition->get_size(),
+            partition->get_range_bounds(),


Shouldn't there be a check whether these pointers are on device executors ?

I think this is a bit of a temporary solution until we fully introduce device-side views for all of our (relevant) types. So I would leave handling that issue until then.

reference/test/distributed/index_map_kernels.cpp

pratikvn · 2025-01-22T12:52:22Z

reference/test/distributed/index_map_kernels.cpp

+TEST_F(IndexMap, CanGetGlobalWithLocalISWithInvalid)
+{
+    gko::array<global_index_type> global_ids(ref);
+    gko::array<local_index_type> local_ids(ref, {5, 4, 10, 3, 2, 1, 0, 100, 4});


The partition contains 18 elements, so I think it would be relevant to check so that global_id=17 is output correctly.

The global_id=17 is not part of the local index set of rank 1, which is used here. I changed it to rank 0 for this test, and now it also outputs gid=17.

reference/test/distributed/partition_kernels.cpp

test/distributed/index_map_kernels.cpp

test/distributed/partition_kernels.cpp

pratikvn

A couple more comments.

core/distributed/index_map_kernels.hpp

core/test/utils/assertions.hpp

test/distributed/index_map_kernels.cpp

yhmtsai

the new testing method GKO_ASSERT_SEGMENTED_ARRAY_EQ needs a test in core/test/utils/assertions_test.cpp as @pratikvn mentioned
some comments from @pratikvn seems not to be checked yet

yhmtsai · 2025-02-13T13:47:56Z

core/distributed/partition.cpp

+    exec->run(partition::make_build_ranges_by_part(
+        part_ids_.get_const_data(), get_num_ranges(), get_num_parts(),
+        range_ids, num_ranges_per_part));
+    ranges_by_part_ = segmented_array<size_type>::create_from_sizes(
+        std::move(range_ids), num_ranges_per_part);


It can be future improvement not in this pr

yhmtsai · 2025-02-13T13:50:24Z

core/distributed/device_partition.hpp

+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#pragma once


I still prefer the official supporting macro approach not relying on the compiler support.

There are only very obscure edge cases where pragma once doesn't work properly (because some filesystems don't make it easy to see if two files have the same identity), I think we could safely move from header guards to #pragma once if we wanted to. Every relevant compiler supports it.

test/distributed/index_map_kernels.cpp

yhmtsai · 2025-02-13T14:18:25Z

reference/distributed/partition_helpers.hpp

+    assert(range_id <
+           std::distance(partition.offsets_begin, partition.offsets_end) - 1);


Suggested change

assert(range_id <

std::distance(partition.offsets_begin, partition.offsets_end) - 1);

assert(range_id <

std::distance(partition.offsets_begin, partition.offsets_end));

do you need -1?

Yes, the offsets array has #range-ids + 1 entries, since it's an offset array. For an empty partition, offsets = {0}, so the begin and end pointers are different.

reference/test/distributed/index_map_kernels.cpp

yhmtsai · 2025-02-13T14:42:54Z

common/cuda_hip/distributed/index_map_kernels.cpp

+        auto local_range_id =
+            it != local_ranges_size ? it : max(int64(0), it - 1);


will this happen? it seems to be broken if we reach the else branch

I realized that I implemented this a bit incorrectly. I'm now using the correct predicate for the binary search, which makes the check here unnecessary. If also added a comment to clarify this a bit.

yhmtsai · 2025-02-13T14:46:21Z

common/cuda_hip/distributed/partition_kernels.cpp

+    thrust::stable_sort_by_key(policy, range_parts_ptr,
+                               range_parts_ptr + num_ranges, range_ids_ptr);


it uses the truth that range ids is also sorted, right?
such that it is sorted by (part_id, range_id)

Yes, although the part_id is not used after this anymore. It's only relevant that the range_id is first sorted by part_id, and then by range_id.

upsj

Some minor change requests and comments

upsj · 2025-02-14T08:10:11Z

test/distributed/partition_kernels.cpp

@@ -69,6 +69,9 @@ class Partition : public CommonTestFixture {
            gko::make_array_view(
                this->exec, dpart->get_num_parts(),
                const_cast<local_index_type*>(dpart->get_part_sizes())));
+
+        GKO_ASSERT_SEGMENTED_ARRAY_EQ(part->get_ranges_by_part(),


Side-node: Since this is a function returning void, this will not prevent subsequent operations from executing.

This should be implemented the same way as the GKO_ASSERT_ARRAY_EQ macro. The segmented_array_equal function will return something, so I don't know if there is anything more to do here.

I mean the overall approach of implementing a custom assert_something function may cause tests to continue executing after the assertion fails, which may or may not be desirable. It's fine here, but something to keep in mind. To fix this, we would need to return testing::AssertionResult and wrap the call in a macro. But nothing to hold back the code here, just a comment that people should be aware of.

reference/distributed/index_map_kernels.cpp

omp/distributed/index_map_kernels.cpp

dpcpp/distributed/partition_kernels.dp.cpp

core/test/utils/assertions.hpp

core/distributed/index_map_kernels.hpp

core/distributed/device_partition.hpp

upsj · 2025-02-14T08:30:59Z

core/distributed/device_partition.hpp

+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#pragma once


There are only very obscure edge cases where pragma once doesn't work properly (because some filesystems don't make it easy to see if two files have the same identity), I think we could safely move from header guards to #pragma once if we wanted to. Every relevant compiler supports it.

yhmtsai · 2025-02-14T10:20:21Z

omp/distributed/index_map_kernels.cpp

-
-

should be kept

yhmtsai · 2025-02-14T10:21:05Z

reference/distributed/index_map_kernels.cpp

-
-
 template <typename LocalIndexType, typename GlobalIndexType>


should kept

pratikvn

LGTM!

upsj

LGTM!

Signed-off-by: Marcel Koch <[email protected]>

Co-authored-by: Pratik Nayak <[email protected]>

- reduce tests - update docs - minor refactoring - fix binary search usage in cuda/hip - refactor - add tests for segmented array assertion Co-authored-by: Yu-Hsiang M. Tsai <[email protected]> Co-authored-by: Pratik Nayak <[email protected]> Co-authored-by: Tobias Ribizel <[email protected]>

sonarqubecloud · 2025-02-28T01:17:49Z

Quality Gate passed

Issues
5 New issues
0 Accepted issues

Measures
0 Security Hotspots
95.7% Coverage on New Code
14.8% Duplication on New Code

See analysis details on SonarQube Cloud

codecov · 2025-02-28T01:23:41Z

Codecov Report

Attention: Patch coverage is 96.53846% with 9 lines in your changes missing coverage. Please review.

Project coverage is 88.58%. Comparing base (d7d3d40) to head (71ac6a4).
Report is 10 commits behind head on develop.

Files with missing lines	Patch %	Lines
core/distributed/index_map.cpp	0.00%	7 Missing ⚠️
core/device_hooks/common_kernels.inc.cpp	0.00%	2 Missing ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           develop    #1707      +/-   ##
===========================================
- Coverage    89.17%   88.58%   -0.59%     
===========================================
  Files          817      818       +1     
  Lines        67816    68151     +335     
===========================================
- Hits         60474    60374     -100     
- Misses        7342     7777     +435

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

MarcelKoch added the 1:ST:ready-for-review This PR is ready for review label Oct 23, 2024

MarcelKoch added this to the Ginkgo 1.9.0 milestone Oct 23, 2024

MarcelKoch self-assigned this Oct 23, 2024

ginkgo-bot added reg:testing This is related to testing. mod:all This touches all Ginkgo modules. labels Oct 23, 2024

MarcelKoch force-pushed the global-to-local branch from 1070d82 to 8392d07 Compare October 23, 2024 15:13

MarcelKoch mentioned this pull request Oct 23, 2024

Use Index map in PGM #1639

Merged

2 tasks

fritzgoebel self-requested a review October 30, 2024 10:36

fritzgoebel mentioned this pull request Oct 31, 2024

Add a Domain decomposition Matrix format #1719

Open

3 tasks

MarcelKoch modified the milestones: Ginkgo 1.9.0, Ginkgo 1.10.0 Dec 9, 2024

MarcelKoch force-pushed the global-to-local branch 2 times, most recently from c43c48a to c3b27c2 Compare December 17, 2024 11:26

MarcelKoch mentioned this pull request Dec 17, 2024

Add specific communicator for neighborhood communication #1588

Merged

9 tasks

yhmtsai reviewed Dec 18, 2024

View reviewed changes

MarcelKoch force-pushed the global-to-local branch from c3b27c2 to bc75a0b Compare January 8, 2025 15:49

pratikvn reviewed Jan 22, 2025

View reviewed changes

MarcelKoch force-pushed the global-to-local branch from 5233e98 to 7863f01 Compare January 24, 2025 11:04

MarcelKoch requested review from pratikvn and yhmtsai January 24, 2025 12:32

pratikvn approved these changes Jan 27, 2025

View reviewed changes

core/distributed/index_map_kernels.hpp Show resolved Hide resolved

core/test/utils/assertions.hpp Show resolved Hide resolved

test/distributed/index_map_kernels.cpp Show resolved Hide resolved

upsj self-requested a review February 13, 2025 11:50

yhmtsai requested changes Feb 13, 2025

View reviewed changes

upsj requested changes Feb 14, 2025

View reviewed changes

MarcelKoch requested a review from upsj February 14, 2025 09:56

MarcelKoch requested review from pratikvn and yhmtsai February 14, 2025 09:56

yhmtsai approved these changes Feb 14, 2025

View reviewed changes

MarcelKoch force-pushed the global-to-local branch from 92b9692 to a75e164 Compare February 14, 2025 10:43

pratikvn approved these changes Feb 14, 2025

View reviewed changes

MarcelKoch force-pushed the global-to-local branch from a75e164 to 7e19ec3 Compare February 14, 2025 12:43

MarcelKoch mentioned this pull request Feb 14, 2025

Adds distributed row gatherer #1589

Open

4 tasks

upsj approved these changes Feb 18, 2025

View reviewed changes

MarcelKoch added 1:ST:ready-to-merge This PR is ready to merge. and removed 1:ST:ready-for-review This PR is ready for review labels Feb 18, 2025

MarcelKoch force-pushed the global-to-local branch 2 times, most recently from 8db4eea to 5b5ddac Compare February 18, 2025 15:46

MarcelKoch changed the base branch from develop to neighborhood-communicator February 20, 2025 09:35

MarcelKoch force-pushed the global-to-local branch from 5b5ddac to 5be2c7f Compare February 20, 2025 09:46

MarcelKoch force-pushed the neighborhood-communicator branch from 37f5bba to d7d3d40 Compare February 25, 2025 07:35

Base automatically changed from neighborhood-communicator to develop February 26, 2025 12:16

[core] fix get_segment for const context

ddc4120

Signed-off-by: Marcel Koch <[email protected]>

MarcelKoch force-pushed the global-to-local branch from 5be2c7f to 285ff97 Compare February 26, 2025 12:31

MarcelKoch and others added 7 commits February 27, 2025 13:37

[dist] provide range-ids segmented by part-id

bf83b1e

Signed-off-by: Marcel Koch <[email protected]>

[dist] add device partition

87380eb

Signed-off-by: Marcel Koch <[email protected]>

[imap] add local to global mapping

447147c

Signed-off-by: Marcel Koch <[email protected]>

[dist] fix to global mapping

bf51756

[dist] add round trip index mapping test

a8531ed

[test] add assertion macro for segmented arrays

9588885

Co-authored-by: Pratik Nayak <[email protected]>

MarcelKoch force-pushed the global-to-local branch from 285ff97 to 71ac6a4 Compare February 27, 2025 12:37

MarcelKoch merged commit 7404a82 into develop Feb 28, 2025
13 of 14 checks passed

MarcelKoch deleted the global-to-local branch February 28, 2025 07:52

	#ifndef GINKGO_PARTITION_HPP
	#ifndef GKO_CORE_DISTRIBUTED_DEVICE_PARTITION_HPP_

		assert(range_id <
		std::distance(partition.offsets_begin, partition.offsets_end) - 1);

		auto local_range_id =
		it != local_ranges_size ? it : max(int64(0), it - 1);

		thrust::stable_sort_by_key(policy, range_parts_ptr,
		range_parts_ptr + num_ranges, range_ids_ptr);



		template <typename LocalIndexType, typename GlobalIndexType>

Add local to global index mapping #1707

Add local to global index mapping #1707

Conversation

MarcelKoch commented Oct 23, 2024

yhmtsai 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

MarcelKoch commented Dec 18, 2024

yhmtsai commented Dec 18, 2024

MarcelKoch commented Dec 18, 2024

yhmtsai commented Dec 18, 2024

MarcelKoch commented Dec 18, 2024

pratikvn 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

pratikvn left a comment

Choose a reason for hiding this comment

yhmtsai 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

upsj 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

pratikvn left a comment

Choose a reason for hiding this comment

upsj left a comment

Choose a reason for hiding this comment

sonarqubecloud bot commented Feb 28, 2025

Quality Gate passed

codecov bot commented Feb 28, 2025 • edited Loading

Codecov Report

codecov bot commented Feb 28, 2025 •

edited

Loading