From db6fc1a2a201354c5dcf43462fd64e35cea2bf22 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Thomas=20Gr=C3=BCtzmacher?= Date: Wed, 10 Jul 2024 17:52:20 +0200 Subject: [PATCH 1/4] Add additional tests for the communicator group --- cuda/test/components/cooperative_groups.cu | 37 ++++++++++++++++ .../test/components/cooperative_groups.dp.cpp | 42 +++++++++++++++++++ .../components/cooperative_groups.hip.cpp | 37 ++++++++++++++++ 3 files changed, 116 insertions(+) diff --git a/cuda/test/components/cooperative_groups.cu b/cuda/test/components/cooperative_groups.cu index df3cef86bb8..077b0121fbd 100644 --- a/cuda/test/components/cooperative_groups.cu +++ b/cuda/test/components/cooperative_groups.cu @@ -223,4 +223,41 @@ TEST_F(CooperativeGroups, SubwarpBallot) { test(cg_subwarp_ballot); } TEST_F(CooperativeGroups, SubwarpBallot2) { test_subwarp(cg_subwarp_ballot); } +__global__ void cg_communicator_categorization(bool*) +{ + auto this_block = group::this_thread_block(); + auto tiled_partition = + group::tiled_partition(this_block); + auto subwarp_partition = group::tiled_partition(this_block); + + using not_group = int; + using this_block_t = decltype(this_block); + using tiled_partition_t = decltype(tiled_partition); + using subwarp_partition_t = decltype(subwarp_partition); + + static_assert(!group::is_group::value && + group::is_group::value && + group::is_group::value && + group::is_group::value, + "Group check doesn't work."); + static_assert( + !group::is_synchronizable_group::value && + group::is_synchronizable_group::value && + group::is_synchronizable_group::value && + group::is_synchronizable_group::value, + "Synchronizable group check doesn't work."); + static_assert( + !group::is_communicator_group::value && + !group::is_communicator_group::value && + group::is_communicator_group::value && + group::is_communicator_group::value, + "Communicator group check doesn't work."); +} + +TEST_F(CooperativeGroups, CorrectCategorization) +{ + test(cg_communicator_categorization); +} + + } // namespace diff --git a/dpcpp/test/components/cooperative_groups.dp.cpp b/dpcpp/test/components/cooperative_groups.dp.cpp index 27e14b62d2d..8667a85713e 100644 --- a/dpcpp/test/components/cooperative_groups.dp.cpp +++ b/dpcpp/test/components/cooperative_groups.dp.cpp @@ -198,6 +198,48 @@ GKO_ENABLE_DEFAULT_CONFIG_CALL(cg_ballot_call, cg_ballot, default_config_list) TEST_P(CooperativeGroups, Ballot) { test_all_subgroup(cg_ballot_call); } +template +void cg_communicator_categorization(bool* s, sycl::nd_item<3> item_ct1) +{ + auto this_block = group::this_thread_block(item_ct1); + auto tiled_partition = + group::tiled_partition(this_block); + + using not_group = int; + using this_block_t = decltype(this_block); + using tiled_partition_t = decltype(tiled_partition); + + static_assert(!group::is_group::value && + group::is_group::value && + group::is_group::value, + "Group check doesn't work."); + static_assert( + !group::is_synchronizable_group::value && + group::is_synchronizable_group::value && + group::is_synchronizable_group::value, + "Synchronizable group check doesn't work."); + static_assert( + !group::is_communicator_group::value && + !group::is_communicator_group::value && + group::is_communicator_group::value, + "Communicator group check doesn't work."); + s[this_block.thread_rank()] = true; +} + +GKO_ENABLE_DEFAULT_HOST_CONFIG_TYPE(cg_communicator_categorization, + cg_communicator_categorization) +GKO_ENABLE_IMPLEMENTATION_CONFIG_SELECTION_TOTYPE( + cg_communicator_categorization, cg_communicator_categorization, DCFG_1D) +GKO_ENABLE_DEFAULT_CONFIG_CALL(cg_communicator_categorization_call, + cg_communicator_categorization, + default_config_list) + +TEST_P(CooperativeGroups, CorrectCategorization) +{ + test_all_subgroup(cg_communicator_categorization_call); +} + + INSTANTIATE_TEST_SUITE_P(DifferentSubgroup, CooperativeGroups, testing::Values(4, 8, 16, 32, 64), testing::PrintToStringParamName()); diff --git a/hip/test/components/cooperative_groups.hip.cpp b/hip/test/components/cooperative_groups.hip.cpp index 06a104a8879..0f71550139c 100644 --- a/hip/test/components/cooperative_groups.hip.cpp +++ b/hip/test/components/cooperative_groups.hip.cpp @@ -242,6 +242,43 @@ TEST_F(CooperativeGroups, SubwarpBallot) { test(cg_subwarp_ballot); } TEST_F(CooperativeGroups, SubwarpBallot2) { test_subwarp(cg_subwarp_ballot); } +__global__ void cg_communicator_categorization(bool*) +{ + auto this_block = group::this_thread_block(); + auto tiled_partition = + group::tiled_partition(this_block); + auto subwarp_partition = group::tiled_partition(this_block); + + using not_group = int; + using this_block_t = decltype(this_block); + using tiled_partition_t = decltype(tiled_partition); + using subwarp_partition_t = decltype(subwarp_partition); + + static_assert(!group::is_group::value && + group::is_group::value && + group::is_group::value && + group::is_group::value, + "Group check doesn't work."); + static_assert( + !group::is_synchronizable_group::value && + group::is_synchronizable_group::value && + group::is_synchronizable_group::value && + group::is_synchronizable_group::value, + "Synchronizable group check doesn't work."); + static_assert( + !group::is_communicator_group::value && + !group::is_communicator_group::value && + group::is_communicator_group::value && + group::is_communicator_group::value, + "Communicator group check doesn't work."); +} + +TEST_F(CooperativeGroups, CorrectCategorization) +{ + test(cg_communicator_categorization); +} + + template __global__ void cg_shuffle_sum(const int num, ValueType* __restrict__ value) { From dc951c356bccb4c863d7eafa8653a395d69878eb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Thomas=20Gr=C3=BCtzmacher?= Date: Thu, 11 Jul 2024 14:13:54 +0200 Subject: [PATCH 2/4] Fix visibility of HIP specialization --- .../test/components/cooperative_groups.dp.cpp | 21 ++++++++++--------- hip/components/cooperative_groups.hip.hpp | 6 +++--- 2 files changed, 14 insertions(+), 13 deletions(-) diff --git a/dpcpp/test/components/cooperative_groups.dp.cpp b/dpcpp/test/components/cooperative_groups.dp.cpp index 8667a85713e..eadd99a6ac5 100644 --- a/dpcpp/test/components/cooperative_groups.dp.cpp +++ b/dpcpp/test/components/cooperative_groups.dp.cpp @@ -213,17 +213,18 @@ void cg_communicator_categorization(bool* s, sycl::nd_item<3> item_ct1) group::is_group::value && group::is_group::value, "Group check doesn't work."); - static_assert( - !group::is_synchronizable_group::value && - group::is_synchronizable_group::value && - group::is_synchronizable_group::value, - "Synchronizable group check doesn't work."); - static_assert( - !group::is_communicator_group::value && - !group::is_communicator_group::value && - group::is_communicator_group::value, - "Communicator group check doesn't work."); + static_assert(!group::is_synchronizable_group::value && + group::is_synchronizable_group::value && + group::is_synchronizable_group::value, + "Synchronizable group check doesn't work."); + static_assert(!group::is_communicator_group::value && + !group::is_communicator_group::value && + group::is_communicator_group::value, + "Communicator group check doesn't work."); + // Make it work with the test framework, which performs 3 tests s[this_block.thread_rank()] = true; + s[this_block.thread_rank() + cfg::subgroup_size] = true; + s[this_block.thread_rank() + 2 * cfg::subgroup_size] = true; } GKO_ENABLE_DEFAULT_HOST_CONFIG_TYPE(cg_communicator_categorization, diff --git a/hip/components/cooperative_groups.hip.hpp b/hip/components/cooperative_groups.hip.hpp index d3dbc44a5c8..11581db0b0c 100644 --- a/hip/components/cooperative_groups.hip.hpp +++ b/hip/components/cooperative_groups.hip.hpp @@ -370,12 +370,12 @@ namespace detail { template -struct is_group_impl> : std::true_type {}; +struct is_group_impl> : std::true_type {}; template -struct is_synchronizable_group_impl> : std::true_type { +struct is_synchronizable_group_impl> : std::true_type { }; template -struct is_communicator_group_impl> : std::true_type {}; +struct is_communicator_group_impl> : std::true_type {}; } // namespace detail From e39ccdb3e4cc07167ebd150714628b7f46af3de5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Thomas=20Gr=C3=BCtzmacher?= Date: Wed, 10 Jul 2024 17:54:53 +0200 Subject: [PATCH 3/4] Fix the communicator group categorization --- cuda/components/cooperative_groups.cuh | 2 +- dpcpp/components/cooperative_groups.dp.hpp | 2 +- hip/components/cooperative_groups.hip.hpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cuda/components/cooperative_groups.cuh b/cuda/components/cooperative_groups.cuh index c4ceca9e409..983ec32f9ac 100644 --- a/cuda/components/cooperative_groups.cuh +++ b/cuda/components/cooperative_groups.cuh @@ -113,7 +113,7 @@ struct is_synchronizable_group_impl : std::false_type {}; template -struct is_communicator_group_impl : std::true_type {}; +struct is_communicator_group_impl : std::false_type {}; } // namespace detail diff --git a/dpcpp/components/cooperative_groups.dp.hpp b/dpcpp/components/cooperative_groups.dp.hpp index c758cf42710..33a107ef3f5 100644 --- a/dpcpp/components/cooperative_groups.dp.hpp +++ b/dpcpp/components/cooperative_groups.dp.hpp @@ -101,7 +101,7 @@ struct is_synchronizable_group_impl : std::false_type {}; template -struct is_communicator_group_impl : std::true_type {}; +struct is_communicator_group_impl : std::false_type {}; } // namespace detail diff --git a/hip/components/cooperative_groups.hip.hpp b/hip/components/cooperative_groups.hip.hpp index 11581db0b0c..2e5d7c0abff 100644 --- a/hip/components/cooperative_groups.hip.hpp +++ b/hip/components/cooperative_groups.hip.hpp @@ -101,7 +101,7 @@ struct is_synchronizable_group_impl : std::false_type {}; template -struct is_communicator_group_impl : std::true_type {}; +struct is_communicator_group_impl : std::false_type {}; } // namespace detail From 14f0e24bec1203f13d2c7ccf291b4153bb5315a9 Mon Sep 17 00:00:00 2001 From: ginkgo-bot Date: Thu, 11 Jul 2024 12:28:28 +0000 Subject: [PATCH 4/4] Format files MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Thomas Grützmacher --- cuda/test/components/cooperative_groups.cu | 11 +++++------ hip/components/cooperative_groups.hip.hpp | 7 ++++--- hip/test/components/cooperative_groups.hip.cpp | 11 +++++------ 3 files changed, 14 insertions(+), 15 deletions(-) diff --git a/cuda/test/components/cooperative_groups.cu b/cuda/test/components/cooperative_groups.cu index 077b0121fbd..0b384cd704e 100644 --- a/cuda/test/components/cooperative_groups.cu +++ b/cuda/test/components/cooperative_groups.cu @@ -246,12 +246,11 @@ __global__ void cg_communicator_categorization(bool*) group::is_synchronizable_group::value && group::is_synchronizable_group::value, "Synchronizable group check doesn't work."); - static_assert( - !group::is_communicator_group::value && - !group::is_communicator_group::value && - group::is_communicator_group::value && - group::is_communicator_group::value, - "Communicator group check doesn't work."); + static_assert(!group::is_communicator_group::value && + !group::is_communicator_group::value && + group::is_communicator_group::value && + group::is_communicator_group::value, + "Communicator group check doesn't work."); } TEST_F(CooperativeGroups, CorrectCategorization) diff --git a/hip/components/cooperative_groups.hip.hpp b/hip/components/cooperative_groups.hip.hpp index 2e5d7c0abff..36618bb7f3e 100644 --- a/hip/components/cooperative_groups.hip.hpp +++ b/hip/components/cooperative_groups.hip.hpp @@ -372,10 +372,11 @@ namespace detail { template struct is_group_impl> : std::true_type {}; template -struct is_synchronizable_group_impl> : std::true_type { -}; +struct is_synchronizable_group_impl> + : std::true_type {}; template -struct is_communicator_group_impl> : std::true_type {}; +struct is_communicator_group_impl> + : std::true_type {}; } // namespace detail diff --git a/hip/test/components/cooperative_groups.hip.cpp b/hip/test/components/cooperative_groups.hip.cpp index 0f71550139c..bd8c79b9849 100644 --- a/hip/test/components/cooperative_groups.hip.cpp +++ b/hip/test/components/cooperative_groups.hip.cpp @@ -265,12 +265,11 @@ __global__ void cg_communicator_categorization(bool*) group::is_synchronizable_group::value && group::is_synchronizable_group::value, "Synchronizable group check doesn't work."); - static_assert( - !group::is_communicator_group::value && - !group::is_communicator_group::value && - group::is_communicator_group::value && - group::is_communicator_group::value, - "Communicator group check doesn't work."); + static_assert(!group::is_communicator_group::value && + !group::is_communicator_group::value && + group::is_communicator_group::value && + group::is_communicator_group::value, + "Communicator group check doesn't work."); } TEST_F(CooperativeGroups, CorrectCategorization)