Skip to content

Commit

Permalink
try config selection in cooperative group test
Browse files Browse the repository at this point in the history
  • Loading branch information
@yhmtsai
yhmtsai committed May 3, 2021
1 parent 5c075fe commit 4540556
Show file tree
Hide file tree
Showing 7 changed files with 95 additions and 25 deletions.
1 change: 1 addition & 0 deletions cmake/create_test.cmake
View file
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@ function(ginkgo_create_dpcpp_test test_name)
add_executable(${TEST_TARGET_NAME} ${test_name}.dp.cpp)
target_compile_features("${TEST_TARGET_NAME}" PUBLIC cxx_std_17)
target_compile_options("${TEST_TARGET_NAME}" PRIVATE "${GINKGO_DPCPP_FLAGS}")
target_link_options("${TEST_TARGET_NAME}" PRIVATE -fsycl-device-code-split=per_kernel)
if (GINKGO_DPCPP_SINGLE_MODE)
target_compile_definitions("${TEST_TARGET_NAME}" PRIVATE GINKGO_DPCPP_SINGLE_MODE=1)
endif()
Expand Down
33 changes: 33 additions & 0 deletions core/synthesizer/implementation_selection.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -70,6 +70,39 @@ namespace syn {
} \
}

#define GKO_ENABLE_IMPLEMENTATION_CONFIG_SELECTION(_name, _callable) \
template <typename Predicate, bool... BoolArgs, int... IntArgs, \
gko::size_type... SizeTArgs, typename... TArgs, \
typename... InferredArgs> \
inline void _name(::gko::syn::value_list<gko::Config>, Predicate, \
::gko::syn::value_list<bool, BoolArgs...>, \
::gko::syn::value_list<int, IntArgs...>, \
::gko::syn::value_list<gko::size_type, SizeTArgs...>, \
::gko::syn::type_list<TArgs...>, InferredArgs...) \
GKO_KERNEL_NOT_FOUND; \
\
template <gko::Config K, gko::Config... Rest, typename Predicate, \
bool... BoolArgs, int... IntArgs, gko::size_type... SizeTArgs, \
typename... TArgs, typename... InferredArgs> \
inline void _name( \
::gko::syn::value_list<gko::Config, K, Rest...>, \
Predicate is_eligible, \
::gko::syn::value_list<bool, BoolArgs...> bool_args, \
::gko::syn::value_list<int, IntArgs...> int_args, \
::gko::syn::value_list<gko::size_type, SizeTArgs...> size_args, \
::gko::syn::type_list<TArgs...> type_args, InferredArgs... args) \
{ \
if (is_eligible(K)) { \
std::cout << "call " << K << std::endl; \
_callable<BoolArgs..., IntArgs..., SizeTArgs..., TArgs..., K>( \
std::forward<InferredArgs>(args)...); \
} else { \
_name(::gko::syn::value_list<gko::Config, Rest...>(), is_eligible, \
bool_args, int_args, size_args, type_args, \
std::forward<InferredArgs>(args)...); \
} \
}


} // namespace syn
} // namespace gko
Expand Down
2 changes: 1 addition & 1 deletion dpcpp/base/dim3.dp.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -54,7 +54,7 @@ sycl::nd_range<3> sycl_nd_range(dim3 grid, dim3 block)
{
auto local_range = block.reverse();
auto global_range = grid.reverse() * local_range;
return sycl::nd_range<3>(global_range, local_range)
return sycl::nd_range<3>(global_range, local_range);
}


Expand Down
9 changes: 5 additions & 4 deletions dpcpp/components/cooperative_groups.dp.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -350,10 +350,11 @@ using detail::thread_block_tile;

// Only support tile_partition with 8, 16, 32.
template <unsigned Size, typename Group>
__dpct_inline__ std::enable_if_t<Size> 1 && Size <= 64 &&
(Size & (Size - 1)),
detail::thread_block_tile<Size>> tiled_partition
[[intel::reqd_sub_group_size(Size)]] (const Group &group)
__dpct_inline__
std::enable_if_t<(Size > 1) && Size <= 64 && !(Size & (Size - 1)),
detail::thread_block_tile<Size>>
tiled_partition
[[intel::reqd_sub_group_size(Size)]] (const Group &group)
{
return detail::thread_block_tile<Size>(group);
}
Expand Down
1 change: 1 addition & 0 deletions dpcpp/test/components/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
@@ -1,3 +1,4 @@
ginkgo_create_test(absolute_array)
ginkgo_create_dpcpp_test(cooperative_groups_kernels)
ginkgo_create_test(fill_array)
ginkgo_create_test(precision_conversion)
70 changes: 50 additions & 20 deletions dpcpp/test/components/cooperative_groups_kernels.dp.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -40,7 +40,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.


#include <gtest/gtest.h>

#include <iostream>

#include <ginkgo/core/base/array.hpp>
#include <ginkgo/core/base/executor.hpp>
Expand Down Expand Up @@ -81,6 +81,19 @@ class CooperativeGroups : public ::testing::Test {
ASSERT_TRUE(success);
}

template <typename Kernel>
void test_all_subgroup(Kernel kernel)
{
auto exec_info = dpcpp->get_const_exec_info();
for (auto &i : exec_info.subgroup_sizes) {
kernel(1, i, 0, dpcpp->get_queue(), dpcpp, dresult.get_data());
result = dresult;
auto success = *result.get_const_data();
ASSERT_TRUE(success);
std::cout << i << " success" << std::endl;
}
}

std::shared_ptr<gko::ReferenceExecutor> ref;
std::shared_ptr<gko::DpcppExecutor> dpcpp;
gko::Array<bool> result;
Expand All @@ -95,20 +108,24 @@ void test_assert(bool *success, bool partial)
}
}

template <Config config>
void cg_shuffle(bool *s, sycl::nd_item<3> item_ct1)
// kernel implementation
template <gko::Config config>
[[intel::reqd_work_group_size(1, 1, gko::get_warp_size(config))]] void
cg_shuffle(bool *s, sycl::nd_item<3> item_ct1)
{
auto group = group::tiled_partition<get_warp_size(config)>(
auto group = group::tiled_partition<gko::get_warp_size(config)>(
group::this_thread_block(item_ct1));
auto i = int(group.thread_rank());
test_assert(s, group.shfl_up(i, 1) == sycl::max(0, (int)(i - 1)));
test_assert(s, group.shfl_down(i, 1) ==
sycl::min((unsigned int)(i + 1),
(unsigned int)(get_warp_size(config) - 1)));
test_assert(s,
group.shfl_down(i, 1) ==
sycl::min((unsigned int)(i + 1),
(unsigned int)(gko::get_warp_size(config) - 1)));
test_assert(s, group.shfl(i, 0) == 0);
}

template <Config config>
// group all kernel things together
template <gko::Config config>
void cg_shuffle_host(dim3 grid, dim3 block, size_t dynamic_shared_memory,
sycl::queue *stream, bool *s)
{
Expand All @@ -120,24 +137,37 @@ void cg_shuffle_host(dim3 grid, dim3 block, size_t dynamic_shared_memory,
});
}

GKO_ENABLE_IMPLEMENTATION_SELECTION(cg_shuffle_config, cg_shuffle_host)
// config selection
GKO_ENABLE_IMPLEMENTATION_CONFIG_SELECTION(cg_shuffle_config, cg_shuffle_host)

void cg_shuffle_config(dim3 grid, dim3 block, size_t dynamic_shared_memory,
sycl::queue *stream,
std::shared_ptr<const DpcppExecutor> exec, bool *s)
// the call
void cg_shuffle_config_call(dim3 grid, dim3 block, size_t dynamic_shared_memory,
sycl::queue *stream,
std::shared_ptr<const gko::DpcppExecutor> exec,
bool *s)
{
auto exec_info = exec->get_exec_info();
auto exec_info = exec->get_const_exec_info();
constexpr auto default_config_list =
::gko::syn::value_list<Config, config_set(32, 32)>();
cg_shuffle_config()(
config_list,
[&exec_info](Config config) { return exec_info.validate(config); },
::gko::syn::value_list<gko::Config, gko::config_set(32, 32),
gko::config_set(16, 16), gko::config_set(8, 8),
gko::config_set(4, 4)>();
std::cout << "block.x " << block.x << std::endl;
cg_shuffle_config(
default_config_list,
// validate
[&exec_info, &block](gko::Config config) {
return exec_info.validate(config) &&
(gko::get_warp_size(config) == block.x);
},
::gko::syn::value_list<bool>(), ::gko::syn::value_list<int>(),
::gko::syn::type_list<>(), grid, block, dynamic_shared_memory, stream,
s);
::gko::syn::value_list<gko::size_type>(), ::gko::syn::type_list<>(),
grid, block, dynamic_shared_memory, stream, s);
}

TEST_F(CooperativeGroups, Shuffle) { test(cg_shuffle_config); }
TEST_F(CooperativeGroups, Shuffle)
{
test_all_subgroup(cg_shuffle_config_call);
}


void cg_all(bool *s, sycl::nd_item<3> item_ct1)
Expand Down
4 changes: 4 additions & 0 deletions include/ginkgo/core/base/executor.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -778,6 +778,10 @@ class Executor : public log::EnableLogging<Executor> {
}
};

public:
const exec_info get_const_exec_info() const { return this->exec_info_; }

protected:
/**
* Gets the exec info struct
*
Expand Down

0 comments on commit 4540556

Please sign in to comment.