Skip to content

Commit 2a3d874

Browse files
stephenswatstephenswat
committed
Add tests for scratch space CCA
This commit adds some tests for running CCA with the scratch space by reducing the thread groups to handle extremely small numbers of cells. This will force the edge case of running out of shared memory to occur.
1 parent 53f5d92 commit 2a3d874

File tree

3 files changed

+117
-64
lines changed

3 files changed

+117
-64
lines changed

‎tests/common/tests/cca_test.hpp

+28
Original file line numberDiff line numberDiff line change
@@ -47,6 +47,17 @@ inline traccc::clustering_config default_ccl_test_config() {
4747
return rv;
4848
}
4949

50+
inline traccc::clustering_config tiny_ccl_test_config() {
51+
traccc::clustering_config rv;
52+
53+
rv.threads_per_partition = 128;
54+
rv.max_cells_per_thread = 1;
55+
rv.target_cells_per_thread = 1;
56+
rv.backup_size_multiplier = 16384;
57+
58+
return rv;
59+
}
60+
5061
class ConnectedComponentAnalysisTests
5162
: public traccc::tests::data_test,
5263
public testing::WithParamInterface<
@@ -95,6 +106,23 @@ class ConnectedComponentAnalysisTests
95106
return out;
96107
}
97108

109+
inline static std::vector<std::string> get_test_files_short(void) {
110+
const std::vector<std::pair<std::string, std::size_t>> cases = {
111+
{"trackml_like", 10},
112+
};
113+
std::vector<std::string> out;
114+
115+
for (const std::pair<std::string, std::size_t> &c : cases) {
116+
for (std::size_t i = 0; i < c.second; ++i) {
117+
std::ostringstream ss;
118+
ss << c.first << "_" << std::setfill('0') << std::setw(10) << i;
119+
out.push_back(ss.str());
120+
}
121+
}
122+
123+
return out;
124+
}
125+
98126
inline void test_connected_component_analysis(ParamType p) {
99127
cca_function_t f = std::get<0>(p);
100128
std::string file_prefix = std::get<1>(p);

‎tests/cuda/test_cca.cpp

+45-32
Original file line numberDiff line numberDiff line change
@@ -13,49 +13,54 @@
1313
#include <vecmem/utils/cuda/async_copy.hpp>
1414

1515
#include "tests/cca_test.hpp"
16+
#include "traccc/clusterization/clustering_config.hpp"
1617
#include "traccc/cuda/clusterization/clusterization_algorithm.hpp"
1718
#include "traccc/cuda/utils/stream.hpp"
1819

1920
namespace {
2021

21-
cca_function_t f = [](const traccc::cell_collection_types::host& cells,
22-
const traccc::cell_module_collection_types::host&
23-
modules) {
24-
std::map<traccc::geometry_id, vecmem::vector<traccc::measurement>> result;
22+
cca_function_t get_f_with(traccc::clustering_config cfg) {
23+
return [cfg](const traccc::cell_collection_types::host& cells,
24+
const traccc::cell_module_collection_types::host& modules) {
25+
std::map<traccc::geometry_id, vecmem::vector<traccc::measurement>>
26+
result;
2527

26-
vecmem::host_memory_resource host_mr;
27-
traccc::cuda::stream stream;
28-
vecmem::cuda::device_memory_resource device_mr;
29-
vecmem::cuda::async_copy copy{stream.cudaStream()};
28+
vecmem::host_memory_resource host_mr;
29+
traccc::cuda::stream stream;
30+
vecmem::cuda::device_memory_resource device_mr;
31+
vecmem::cuda::async_copy copy{stream.cudaStream()};
3032

31-
traccc::cuda::clusterization_algorithm cc({device_mr}, copy, stream,
32-
default_ccl_test_config());
33+
traccc::cuda::clusterization_algorithm cc({device_mr}, copy, stream,
34+
cfg);
3335

34-
traccc::cell_collection_types::buffer cells_buffer{
35-
static_cast<traccc::cell_collection_types::buffer::size_type>(
36-
cells.size()),
37-
device_mr};
38-
copy.setup(cells_buffer);
39-
copy(vecmem::get_data(cells), cells_buffer)->ignore();
36+
traccc::cell_collection_types::buffer cells_buffer{
37+
static_cast<traccc::cell_collection_types::buffer::size_type>(
38+
cells.size()),
39+
device_mr};
40+
copy.setup(cells_buffer);
41+
copy(vecmem::get_data(cells), cells_buffer)->ignore();
4042

41-
traccc::cell_module_collection_types::buffer modules_buffer{
42-
static_cast<traccc::cell_module_collection_types::buffer::size_type>(
43-
modules.size()),
44-
device_mr};
45-
copy.setup(modules_buffer);
46-
copy(vecmem::get_data(modules), modules_buffer)->ignore();
43+
traccc::cell_module_collection_types::buffer modules_buffer{
44+
static_cast<
45+
traccc::cell_module_collection_types::buffer::size_type>(
46+
modules.size()),
47+
device_mr};
48+
copy.setup(modules_buffer);
49+
copy(vecmem::get_data(modules), modules_buffer)->ignore();
4750

48-
auto measurements_buffer = cc(cells_buffer, modules_buffer);
49-
traccc::measurement_collection_types::host measurements{&host_mr};
50-
copy(measurements_buffer, measurements)->wait();
51+
auto measurements_buffer = cc(cells_buffer, modules_buffer);
52+
traccc::measurement_collection_types::host measurements{&host_mr};
53+
copy(measurements_buffer, measurements)->wait();
5154

52-
for (std::size_t i = 0; i < measurements.size(); i++) {
53-
result[modules.at(measurements.at(i).module_link).surface_link.value()]
54-
.push_back(measurements.at(i));
55-
}
55+
for (std::size_t i = 0; i < measurements.size(); i++) {
56+
result[modules.at(measurements.at(i).module_link)
57+
.surface_link.value()]
58+
.push_back(measurements.at(i));
59+
}
5660

57-
return result;
58-
};
61+
return result;
62+
};
63+
}
5964
} // namespace
6065

6166
TEST_P(ConnectedComponentAnalysisTests, Run) {
@@ -65,6 +70,14 @@ TEST_P(ConnectedComponentAnalysisTests, Run) {
6570
INSTANTIATE_TEST_SUITE_P(
6671
CUDAFastSvAlgorithm, ConnectedComponentAnalysisTests,
6772
::testing::Combine(
68-
::testing::Values(f),
73+
::testing::Values(get_f_with(default_ccl_test_config())),
6974
::testing::ValuesIn(ConnectedComponentAnalysisTests::get_test_files())),
7075
ConnectedComponentAnalysisTests::get_test_name);
76+
77+
INSTANTIATE_TEST_SUITE_P(
78+
CUDAFastSvAlgorithmWithScratch, ConnectedComponentAnalysisTests,
79+
::testing::Combine(
80+
::testing::Values(get_f_with(tiny_ccl_test_config())),
81+
::testing::ValuesIn(
82+
ConnectedComponentAnalysisTests::get_test_files_short())),
83+
ConnectedComponentAnalysisTests::get_test_name);

‎tests/sycl/test_cca.sycl

+44-32
Original file line numberDiff line numberDiff line change
@@ -19,44 +19,48 @@
1919

2020
namespace {
2121

22-
cca_function_t f = [](const traccc::cell_collection_types::host& cells,
23-
const traccc::cell_module_collection_types::host&
24-
modules) {
25-
std::map<traccc::geometry_id, vecmem::vector<traccc::measurement>> result;
22+
cca_function_t get_f_with(traccc::clustering_config cfg) {
23+
return [cfg](const traccc::cell_collection_types::host& cells,
24+
const traccc::cell_module_collection_types::host& modules) {
25+
std::map<traccc::geometry_id, vecmem::vector<traccc::measurement>>
26+
result;
2627

27-
vecmem::host_memory_resource host_mr;
28-
cl::sycl::queue queue;
29-
vecmem::sycl::device_memory_resource device_mr;
30-
vecmem::sycl::async_copy copy{&queue};
28+
vecmem::host_memory_resource host_mr;
29+
cl::sycl::queue queue;
30+
vecmem::sycl::device_memory_resource device_mr;
31+
vecmem::sycl::async_copy copy{&queue};
3132

32-
traccc::sycl::clusterization_algorithm cc({device_mr}, copy, &queue,
33-
default_ccl_test_config());
33+
traccc::sycl::clusterization_algorithm cc({device_mr}, copy, &queue,
34+
cfg);
3435

35-
traccc::cell_collection_types::buffer cells_buffer{
36-
static_cast<traccc::cell_collection_types::buffer::size_type>(
37-
cells.size()),
38-
device_mr};
39-
copy.setup(cells_buffer);
40-
copy(vecmem::get_data(cells), cells_buffer)->ignore();
36+
traccc::cell_collection_types::buffer cells_buffer{
37+
static_cast<traccc::cell_collection_types::buffer::size_type>(
38+
cells.size()),
39+
device_mr};
40+
copy.setup(cells_buffer);
41+
copy(vecmem::get_data(cells), cells_buffer)->ignore();
4142

42-
traccc::cell_module_collection_types::buffer modules_buffer{
43-
static_cast<traccc::cell_module_collection_types::buffer::size_type>(
44-
modules.size()),
45-
device_mr};
46-
copy.setup(modules_buffer);
47-
copy(vecmem::get_data(modules), modules_buffer)->ignore();
43+
traccc::cell_module_collection_types::buffer modules_buffer{
44+
static_cast<
45+
traccc::cell_module_collection_types::buffer::size_type>(
46+
modules.size()),
47+
device_mr};
48+
copy.setup(modules_buffer);
49+
copy(vecmem::get_data(modules), modules_buffer)->ignore();
4850

49-
auto measurements_buffer = cc(cells_buffer, modules_buffer);
50-
traccc::measurement_collection_types::host measurements{&host_mr};
51-
copy(measurements_buffer, measurements)->wait();
51+
auto measurements_buffer = cc(cells_buffer, modules_buffer);
52+
traccc::measurement_collection_types::host measurements{&host_mr};
53+
copy(measurements_buffer, measurements)->wait();
5254

53-
for (std::size_t i = 0; i < measurements.size(); i++) {
54-
result[modules.at(measurements.at(i).module_link).surface_link.value()]
55-
.push_back(measurements.at(i));
56-
}
55+
for (std::size_t i = 0; i < measurements.size(); i++) {
56+
result[modules.at(measurements.at(i).module_link)
57+
.surface_link.value()]
58+
.push_back(measurements.at(i));
59+
}
5760

58-
return result;
59-
};
61+
return result;
62+
};
63+
}
6064
} // namespace
6165

6266
TEST_P(ConnectedComponentAnalysisTests, Run) {
@@ -66,6 +70,14 @@ TEST_P(ConnectedComponentAnalysisTests, Run) {
6670
INSTANTIATE_TEST_SUITE_P(
6771
SYCLFastSvAlgorithm, ConnectedComponentAnalysisTests,
6872
::testing::Combine(
69-
::testing::Values(f),
73+
::testing::Values(get_f_with(default_ccl_test_config())),
7074
::testing::ValuesIn(ConnectedComponentAnalysisTests::get_test_files())),
7175
ConnectedComponentAnalysisTests::get_test_name);
76+
77+
INSTANTIATE_TEST_SUITE_P(
78+
SYCLFastSvAlgorithmWithScratch, ConnectedComponentAnalysisTests,
79+
::testing::Combine(
80+
::testing::Values(get_f_with(tiny_ccl_test_config())),
81+
::testing::ValuesIn(
82+
ConnectedComponentAnalysisTests::get_test_files_short())),
83+
ConnectedComponentAnalysisTests::get_test_name);

0 commit comments

Comments
 (0)