diff --git a/DataFormats/SoATemplate/README.md b/DataFormats/SoATemplate/README.md index 81bf9d859019e..63c356fa83c76 100644 --- a/DataFormats/SoATemplate/README.md +++ b/DataFormats/SoATemplate/README.md @@ -101,6 +101,11 @@ In addition to those fully parametrized templates, two further levels of paramet mirroring the structure of the underlying structs. The blocks are built via composition, and access to individual layouts and views is provided by name. +`SoABlocks` also have the possibility of generating methods for the `View` and `ConstView` classes using the macros `SOA_VIEW_METHODS` +and `SOA_CONST_VIEW_METHODS`. Like the macros for the element methods, this can also be called only once, and if more methods +have to be generated, they must be listed inside the same macro call. Since these methods can be called from the device, +they must be prefixed with the `SOA_HOST_DEVICE` macro and, when possible, with the `constexpr` keyword. + TODOs: - Add introspection utilities to print the structure and layout of a `SoABlocks` instance. @@ -234,6 +239,43 @@ using SoAView = SoA::View; using SoAConstView = SoA::ConstView; ``` +as well as methods that operate on the View: + +```C++ +GENERATE_SOA_LAYOUT(PositionLayout, + SOA_COLUMN(float, x), + SOA_COLUMN(float, y), + SOA_COLUMN(float, z)) + +GENERATE_SOA_LAYOUT(VelocityLayout, + SOA_COLUMN(float, vx), + SOA_COLUMN(float, vy), + SOA_COLUMN(float, vz)) + +GENERATE_SOA_BLOCKS(PointsLayout, + SOA_BLOCK(position, PositionLayout), + SOA_BLOCK(velocity, VelocityLayout), + SOA_VIEW_METHODS( + SOA_HOST_DEVICE void update_position(uint32_t i, float time) { + auto pos = this->position()[i]; + auto vel = this->velocity()[i]; + pos.x() += vel.vx() * time; + pos.y() += vel.vy() * time; + pos.z() += vel.vz() * time; + } + ), + SOA_CONST_VIEW_METHODS( + SOA_HOST_DEVICE auto distance2(uint32_t i, uint32_t j) const { + auto pi = this->position()[i]; + auto pj = this->position()[j]; + return (pi.x() - pj.x()) * (pi.x() - pj.x()) + + (pi.y() - pj.y()) * (pi.y() - pj.y()) + + (pi.z() - pj.z()) * (pi.z() - pj.z()); + } + ) +) +``` + The buffer of the proper size is allocated, and the layout is populated with: ```C++ diff --git a/DataFormats/SoATemplate/interface/SoABlocks.h b/DataFormats/SoATemplate/interface/SoABlocks.h index b23737641bc74..533ada52074a3 100644 --- a/DataFormats/SoATemplate/interface/SoABlocks.h +++ b/DataFormats/SoATemplate/interface/SoABlocks.h @@ -466,6 +466,7 @@ \ /* Accessors for the const views for each block */ \ _ITERATE_ON_ALL(_DECLARE_ACCESSORS_CONST_VIEW_BLOCKS, ~, __VA_ARGS__) \ + ENUM_IF_VALID(_ITERATE_ON_ALL(GENERATE_CONST_VIEW_METHODS, ~, __VA_ARGS__)) \ \ private: \ _ITERATE_ON_ALL(_DECLARE_MEMBERS_CONST_VIEW_BLOCKS, ~, __VA_ARGS__) \ @@ -548,6 +549,7 @@ \ /* Accessors for the views for each block */ \ _ITERATE_ON_ALL(_DECLARE_ACCESSORS_VIEW_BLOCKS, ~, __VA_ARGS__) \ + ENUM_IF_VALID(_ITERATE_ON_ALL(GENERATE_VIEW_METHODS, ~, __VA_ARGS__)) \ \ /* Data members inherited from the ConstView */ \ }; \ diff --git a/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cc b/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cc index c73cd059384ce..db4a96e87da40 100644 --- a/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cc +++ b/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cc @@ -30,7 +30,7 @@ TEST_CASE("SoACustomizedMethods") { } view.detectorType() = 42; - SECTION("ConstView methods") { + SECTION("ConstElement methods") { // arrays of norms std::array position_norms; std::array velocity_norms; @@ -46,7 +46,7 @@ TEST_CASE("SoACustomizedMethods") { } } - SECTION("View methods") { + SECTION("Element methods") { // array of times std::array times; @@ -71,4 +71,33 @@ TEST_CASE("SoACustomizedMethods") { REQUIRE_THAT(view[i].square_norm_velocity(), Catch::Matchers::WithinAbs(1., 1.e-9)); } } + + const auto points_sizes = std::array{{2, 2}}; + const std::size_t blocksBufferSize = Points::computeDataSize(points_sizes); + + std::unique_ptr points_buffer{ + reinterpret_cast(aligned_alloc(Points::alignment, blocksBufferSize)), std::free}; + + Points points(points_buffer.get(), points_sizes); + PointsView points_view{points}; + PointsConstView points_const_view{points}; + points_view.position()[0].x() = 2.f; + points_view.position()[0].y() = 4.f; + points_view.position()[0].z() = 3.f; + points_view.position()[1].x() = 1.f; + points_view.position()[1].y() = 1.f; + points_view.position()[1].z() = 1.f; + + SECTION("View methods") { REQUIRE(points_const_view.distance2(0, 1) == 14.f); } + + SECTION("ConstView methods") { + points_view.velocity()[0].vx() = 1.f; + points_view.velocity()[0].vy() = 3.f; + points_view.velocity()[0].vz() = 5.f; + const auto time = .5f; + points_view.update_position(0, time); + REQUIRE(points_view.position()[0].x() == 2.5f); + REQUIRE(points_view.position()[0].y() == 5.5f); + REQUIRE(points_view.position()[0].z() == 5.5f); + } } diff --git a/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cu b/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cu index 2941aadc6d059..c6f2a56dd43a5 100644 --- a/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cu +++ b/DataFormats/SoATemplate/test/SoACustomizedMethods_t.cu @@ -24,6 +24,15 @@ __global__ void checkNormalise(SoAView soaView, double* checkTimesFunction) { soaView[i].normalise(); } +__global__ void checkPointsDistance(PointsConstView view, bool* result) { *result &= (view.distance2(0, 1) == 14.f); } + +__global__ void checkPointsPositionUpdate(PointsView view, float time, bool* result) { + view.update_position(0, time); + *result &= (view.position()[0].x() == 2.5f); + *result &= (view.position()[0].y() == 5.5f); + *result &= (view.position()[0].z() == 5.5f); +} + TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") { // common number of elements for the SoAs const std::size_t elems = 10; @@ -69,7 +78,7 @@ TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") { // Host → Device copy CUDA_CHECK(cudaMemcpy(d_buf, h_buf, bufferSize, cudaMemcpyHostToDevice)); - SECTION("ConstView methods CUDA") { + SECTION("ConstElement methods CUDA") { calculateNorm<<<(elems + 255) / 256, 256>>>(d_Constview, d_position_norms, d_velocity_norms); CUDA_CHECK(cudaMemcpy(h_position_norms.data(), d_position_norms, elems * sizeof(float), cudaMemcpyDeviceToHost)); @@ -88,7 +97,7 @@ TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") { } } - SECTION("View methods CUDA") { + SECTION("Element methods CUDA") { std::array times; // Check for the correctness of the time() function @@ -116,10 +125,55 @@ TEST_CASE("SoACustomizedMethods CUDA", "[SoACustomizedMethods][cuda]") { } } + const auto points_sizes = std::array{{2, 2}}; + const std::size_t blocksBufferSize = Points::computeDataSize(points_sizes); + + std::byte* points_buffer_host = nullptr; + CUDA_CHECK(cudaMallocHost(&points_buffer_host, blocksBufferSize)); + + Points h_points(points_buffer_host, points_sizes); + PointsView h_points_view{h_points}; + h_points_view.position()[0].x() = 2.f; + h_points_view.position()[0].y() = 4.f; + h_points_view.position()[0].z() = 3.f; + h_points_view.position()[1].x() = 1.f; + h_points_view.position()[1].y() = 1.f; + h_points_view.position()[1].z() = 1.f; + h_points_view.velocity()[0].vx() = 1.f; + h_points_view.velocity()[0].vy() = 3.f; + h_points_view.velocity()[0].vz() = 5.f; + + std::byte* points_buffer_device = nullptr; + CUDA_CHECK(cudaMalloc(&points_buffer_device, blocksBufferSize)); + + Points d_points(points_buffer_device, points_sizes); + PointsView d_points_view{d_points}; + PointsConstView d_points_const_view{d_points}; + + CUDA_CHECK(cudaMemcpy(points_buffer_device, points_buffer_host, blocksBufferSize, cudaMemcpyHostToDevice)); + + bool* d_result = nullptr; + CUDA_CHECK(cudaMalloc(&d_result, sizeof(bool))); + CUDA_CHECK(cudaMemset(d_result, 1, sizeof(bool))); + + SECTION("View methods") { checkPointsDistance<<<1, 1>>>(d_points_const_view, d_result); } + + SECTION("ConstView methods") { + const auto time = .5f; + checkPointsPositionUpdate<<<1, 1>>>(d_points_view, time, d_result); + } + + bool h_result; + CUDA_CHECK(cudaMemcpy(&h_result, d_result, sizeof(bool), cudaMemcpyDeviceToHost)); + REQUIRE(h_result); + // ===== cleanup ===== CUDA_CHECK(cudaFree(d_position_norms)); CUDA_CHECK(cudaFree(d_velocity_norms)); CUDA_CHECK(cudaFree(d_times)); CUDA_CHECK(cudaFree(d_buf)); + CUDA_CHECK(cudaFree(points_buffer_device)); + CUDA_CHECK(cudaFree(d_result)); CUDA_CHECK(cudaFreeHost(h_buf)); + CUDA_CHECK(cudaFreeHost(points_buffer_host)); } diff --git a/DataFormats/SoATemplate/test/SoACustomizedMethods_t.hip.cc b/DataFormats/SoATemplate/test/SoACustomizedMethods_t.hip.cc index 997efcc8b2173..1dbf32a85197f 100644 --- a/DataFormats/SoATemplate/test/SoACustomizedMethods_t.hip.cc +++ b/DataFormats/SoATemplate/test/SoACustomizedMethods_t.hip.cc @@ -26,6 +26,15 @@ __global__ void checkNormalise(SoAView soaView, double* checkTimesFunction) { soaView[i].normalise(); } +__global__ void checkPointsDistance(PointsConstView view, bool* result) { *result &= (view.distance2(0, 1) == 14.f); } + +__global__ void checkPointsPositionUpdate(PointsView view, float time, bool* result) { + view.update_position(0, time); + *result &= (view.position()[0].x() == 2.5f); + *result &= (view.position()[0].y() == 5.5f); + *result &= (view.position()[0].z() == 5.5f); +} + TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") { // common number of elements for the SoAs const std::size_t elems = 10; @@ -71,7 +80,7 @@ TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") { // Host → Device copy HIP_CHECK(hipMemcpy(d_buf, h_buf, bufferSize, hipMemcpyHostToDevice)); - SECTION("ConstView methods HIP") { + SECTION("ConstElement methods HIP") { calculateNorm<<<(elems + 255) / 256, 256>>>(d_Constview, d_position_norms, d_velocity_norms); HIP_CHECK(hipMemcpy(h_position_norms.data(), d_position_norms, elems * sizeof(float), hipMemcpyDeviceToHost)); @@ -90,7 +99,7 @@ TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") { } } - SECTION("View methods HIP") { + SECTION("Element methods HIP") { std::array times; // Check for the correctness of the time() function @@ -118,10 +127,55 @@ TEST_CASE("SoACustomizedMethods hip", "[SoACustomizedMethods][hip]") { } } + const auto points_sizes = std::array{{2, 2}}; + const std::size_t blocksBufferSize = Points::computeDataSize(points_sizes); + + std::byte* points_buffer_host = nullptr; + HIP_CHECK(hipHostMalloc(&points_buffer_host, blocksBufferSize)); + + Points h_points(points_buffer_host, points_sizes); + PointsView h_points_view{h_points}; + h_points_view.position()[0].x() = 2.f; + h_points_view.position()[0].y() = 4.f; + h_points_view.position()[0].z() = 3.f; + h_points_view.position()[1].x() = 1.f; + h_points_view.position()[1].y() = 1.f; + h_points_view.position()[1].z() = 1.f; + h_points_view.velocity()[0].vx() = 1.f; + h_points_view.velocity()[0].vy() = 3.f; + h_points_view.velocity()[0].vz() = 5.f; + + std::byte* points_buffer_device = nullptr; + HIP_CHECK(hipMalloc(&points_buffer_device, blocksBufferSize)); + + Points d_points(points_buffer_device, points_sizes); + PointsView d_points_view{d_points}; + PointsConstView d_points_const_view{d_points}; + + HIP_CHECK(hipMemcpy(points_buffer_device, points_buffer_host, blocksBufferSize, hipMemcpyHostToDevice)); + + bool* d_result = nullptr; + HIP_CHECK(hipMalloc(&d_result, sizeof(bool))); + HIP_CHECK(hipMemset(d_result, 1, sizeof(bool))); + + SECTION("View methods") { checkPointsDistance<<<1, 1>>>(d_points_const_view, d_result); } + + SECTION("ConstView methods") { + const auto time = .5f; + checkPointsPositionUpdate<<<1, 1>>>(d_points_view, time, d_result); + } + + bool h_result; + HIP_CHECK(hipMemcpy(&h_result, d_result, sizeof(bool), hipMemcpyDeviceToHost)); + REQUIRE(h_result); + // ===== cleanup ===== HIP_CHECK(hipFree(d_position_norms)); HIP_CHECK(hipFree(d_velocity_norms)); HIP_CHECK(hipFree(d_times)); HIP_CHECK(hipFree(d_buf)); + HIP_CHECK(hipFree(points_buffer_device)); + HIP_CHECK(hipFree(d_result)); HIP_CHECK(hipFreeHost(h_buf)); + HIP_CHECK(hipFreeHost(points_buffer_host)); } diff --git a/DataFormats/SoATemplate/test/SoADefinition_CustomizedMethods.h b/DataFormats/SoATemplate/test/SoADefinition_CustomizedMethods.h index 128c26e04c73c..cad4984d89780 100644 --- a/DataFormats/SoATemplate/test/SoADefinition_CustomizedMethods.h +++ b/DataFormats/SoATemplate/test/SoADefinition_CustomizedMethods.h @@ -1,6 +1,7 @@ #include #include "DataFormats/SoATemplate/interface/SoALayout.h" +#include "DataFormats/SoATemplate/interface/SoABlocks.h" GENERATE_SOA_LAYOUT(SoATemplate, SOA_COLUMN(float, x), @@ -18,21 +19,21 @@ GENERATE_SOA_LAYOUT(SoATemplate, x() /= norm_position; y() /= norm_position; z() /= norm_position; - }; + } double norm_velocity = square_norm_velocity(); if (norm_velocity > 0.0f) { v_x() /= norm_velocity; v_y() /= norm_velocity; v_z() /= norm_velocity; - }; + } }), SOA_CONST_ELEMENT_METHODS( SOA_HOST_DEVICE float square_norm_position() - const { return sqrt(x() * x() + y() * y() + z() * z()); }; + const { return sqrt(x() * x() + y() * y() + z() * z()); } SOA_HOST_DEVICE double square_norm_velocity() - const { return sqrt(v_x() * v_x() + v_y() * v_y() + v_z() * v_z()); }; + const { return sqrt(v_x() * v_x() + v_y() * v_y() + v_z() * v_z()); } template SOA_HOST_DEVICE static auto time(T1 pos, T2 vel) { @@ -46,3 +47,42 @@ GENERATE_SOA_LAYOUT(SoATemplate, using SoA = SoATemplate<>; using SoAView = SoA::View; using SoAConstView = SoA::ConstView; + +// clang-format off +GENERATE_SOA_LAYOUT(PositionLayout, + SOA_COLUMN(float, x), + SOA_COLUMN(float, y), + SOA_COLUMN(float, z)) + +GENERATE_SOA_LAYOUT(VelocityLayout, + SOA_COLUMN(float, vx), + SOA_COLUMN(float, vy), + SOA_COLUMN(float, vz)) + +GENERATE_SOA_BLOCKS(PointsLayout, + SOA_BLOCK(position, PositionLayout), + SOA_BLOCK(velocity, VelocityLayout), + SOA_VIEW_METHODS( + SOA_HOST_DEVICE void update_position(uint32_t i, float time) { + auto pos = this->position()[i]; + auto vel = this->velocity()[i]; + pos.x() += vel.vx() * time; + pos.y() += vel.vy() * time; + pos.z() += vel.vz() * time; + } + ), + SOA_CONST_VIEW_METHODS( + SOA_HOST_DEVICE auto distance2(uint32_t i, uint32_t j) const { + auto pi = this->position()[i]; + auto pj = this->position()[j]; + return (pi.x() - pj.x()) * (pi.x() - pj.x()) + + (pi.y() - pj.y()) * (pi.y() - pj.y()) + + (pi.z() - pj.z()) * (pi.z() - pj.z()); + } + ) +) +// clang-format on + +using Points = PointsLayout<>; +using PointsView = Points::View; +using PointsConstView = Points::ConstView; diff --git a/DataFormats/TICL/BuildFile.xml b/DataFormats/TICL/BuildFile.xml new file mode 100644 index 0000000000000..fdf9d63b89a7f --- /dev/null +++ b/DataFormats/TICL/BuildFile.xml @@ -0,0 +1,3 @@ + + + diff --git a/DataFormats/TICL/interface/AssociationMap.h b/DataFormats/TICL/interface/AssociationMap.h new file mode 100644 index 0000000000000..26cc7da2d468a --- /dev/null +++ b/DataFormats/TICL/interface/AssociationMap.h @@ -0,0 +1,82 @@ + +#ifndef DataFormats_TICL_interface_AssociationMap_h +#define DataFormats_TICL_interface_AssociationMap_h + +#include "DataFormats/SoATemplate/interface/SoALayout.h" +#include "DataFormats/SoATemplate/interface/SoABlocks.h" +#include +#include +#include + +namespace ticl { + + namespace concepts { + + template + concept trivially_copyable = std::is_trivially_copyable_v; + + } + + // clang-format off + template + struct AssociationMapLayout { + GENERATE_SOA_LAYOUT(ContentBuffersLayout, SOA_COLUMN(TMapped, values)) + GENERATE_SOA_LAYOUT(OffsetBufferLayout, SOA_COLUMN(TKey, keys_offsets)) + + template + class OffsetsLayout final : public OffsetBufferLayout { + using Parent = OffsetBufferLayout; + + public: + OffsetsLayout() : Parent() {} + OffsetsLayout(std::byte* mem, cms::soa::size_type elements) : Parent(mem, elements + 1) {} + + static constexpr auto computeDataSize(std::size_t size) { + return Parent::computeDataSize(size + 1); + } + }; + + GENERATE_SOA_BLOCKS(Layout, + SOA_BLOCK(content, ContentBuffersLayout), + SOA_BLOCK(offsets, OffsetsLayout), + SOA_VIEW_METHODS( + constexpr SOA_HOST_DEVICE auto operator[](TKey key) { + const auto offset = this->offsets()[key].keys_offsets(); + const auto size = this->count(key); + return std::span{this->content().values().data() + offset, static_cast(size)}; + } + ), + SOA_CONST_VIEW_METHODS( + constexpr SOA_HOST_DEVICE auto operator[](TKey key) const { + const auto offset = this->offsets()[key].keys_offsets(); + const auto size = this->count(key); + return std::span{this->content().values().data() + offset, static_cast(size)}; + } + constexpr SOA_HOST_DEVICE auto contains(TKey key) const { + return this->count(key) > 0; + } + constexpr SOA_HOST_DEVICE auto count(TKey key) const { + return this->offsets()[key + 1].keys_offsets() - this->offsets()[key].keys_offsets(); + } + constexpr SOA_HOST_DEVICE auto keys() const { + return this->offsets().metadata().size() - 1; + } + constexpr SOA_HOST_DEVICE auto size() const { + return this->content().metadata().size(); + } + ) + ) + }; + // clang-format on + + template + using AssociationMap = typename AssociationMapLayout::template Layout<>; + template + using AssociationMapView = typename AssociationMap::View; + template + using AssociationMapConstView = typename AssociationMap::ConstView; + +} // namespace ticl + +#endif diff --git a/DataFormats/TICL/interface/FillAssociator.h b/DataFormats/TICL/interface/FillAssociator.h new file mode 100644 index 0000000000000..fbcb3fb04b5ed --- /dev/null +++ b/DataFormats/TICL/interface/FillAssociator.h @@ -0,0 +1,23 @@ + +#ifndef DataFormats_TICL_interface_FillAssociator_h +#define DataFormats_TICL_interface_FillAssociator_h + +#include "DataFormats/TICL/interface/AssociationMap.h" +#include "DataFormats/TICL/interface/detail/FillAssociator.h" +#include +#include + +namespace ticl::associator { + + template + requires alpaka::isQueue + ALPAKA_FN_HOST auto fill(TQueue& queue, + ticl::AssociationMapView& map, + std::span keys, + std::span values) { + detail::fill(queue, map, keys, values); + } + +} // namespace ticl::associator + +#endif diff --git a/DataFormats/TICL/interface/detail/FillAssociator.h b/DataFormats/TICL/interface/detail/FillAssociator.h new file mode 100644 index 0000000000000..003c47bb0b86d --- /dev/null +++ b/DataFormats/TICL/interface/detail/FillAssociator.h @@ -0,0 +1,88 @@ + +#ifndef DataFormats_TICL_interface_detail_FillAssociator_h +#define DataFormats_TICL_interface_detail_FillAssociator_h + +#include "DataFormats/TICL/interface/AssociationMap.h" +#include "HeterogeneousCore/AlpakaInterface/interface/prefixScan.h" +#include +#include +#include + +namespace ticl::associator::detail { + + struct KernelComputeAssociationSizes { + template + ALPAKA_FN_ACC void operator()(const TAcc& acc, + std::span keys, + TKey* keys_counts, + std::size_t size) const { + for (auto i : alpaka::uniformElements(acc, size)) { + alpaka::atomicAdd(acc, &keys_counts[keys[i]], TKey{1}); + } + } + }; + + struct KernelFillAssociator { + template + ALPAKA_FN_ACC void operator()(const TAcc& acc, + ticl::AssociationMapView view, + std::span keys, + std::span values, + TKey* temp_offsets) const { + for (auto i : alpaka::uniformElements(acc, values.size())) { + const auto key = keys[i]; + const auto offset = alpaka::atomicAdd(acc, &temp_offsets[key], TKey{1}); + view.content().values()[offset] = values[i]; + } + } + }; + + template + requires alpaka::isQueue + ALPAKA_FN_HOST auto fill(TQueue& queue, + ticl::AssociationMapView& map, + std::span keys, + std::span values) { + using namespace ::cms::alpakatools; + + const auto nkeys = map.keys(); + const auto nvalues = map.content().metadata().size(); + + assert(keys.size() == values.size()); + assert(static_cast(values.size()) <= nvalues); + + const auto blocksize = 1024u; + const auto gridsize = divide_up_by(keys.size(), blocksize); + const auto workdiv = make_workdiv(gridsize, blocksize); + auto keys_counts = make_device_buffer(queue, nkeys); + alpaka::memset(queue, keys_counts, 0); + alpaka::exec(queue, workdiv, KernelComputeAssociationSizes{}, keys, keys_counts.data(), keys.size()); + + // prepare for prefix scan + auto block_counter = make_device_buffer(queue); + alpaka::memset(queue, block_counter, 0); + auto temp_offsets = make_device_buffer(queue, nkeys + 1); + alpaka::memset(queue, temp_offsets, 0); + const auto blocksize_multiblockscan = 1024; + auto gridsize_multiblockscan = divide_up_by(nkeys, blocksize_multiblockscan); + const auto workdiv_multiblockscan = make_workdiv(gridsize_multiblockscan, blocksize_multiblockscan); + auto warp_size = alpaka::getPreferredWarpSize(alpaka::getDev(queue)); + alpaka::exec(queue, + workdiv_multiblockscan, + multiBlockPrefixScan{}, + keys_counts.data(), + temp_offsets.data() + 1, + nkeys, + gridsize_multiblockscan, + block_counter.data(), + warp_size); + + alpaka::memcpy(queue, + make_device_view(queue, map.offsets().keys_offsets().data(), nkeys + 1), + make_device_view(queue, temp_offsets.data(), nkeys + 1)); + alpaka::exec(queue, workdiv, KernelFillAssociator{}, map, keys, values, temp_offsets.data()); + } + +} // namespace ticl::associator::detail + +#endif diff --git a/DataFormats/TICL/test/BuildFile.xml b/DataFormats/TICL/test/BuildFile.xml new file mode 100644 index 0000000000000..011735994758f --- /dev/null +++ b/DataFormats/TICL/test/BuildFile.xml @@ -0,0 +1,7 @@ + + + + + + + diff --git a/DataFormats/TICL/test/alpaka/TestAssociationMap.dev.cc b/DataFormats/TICL/test/alpaka/TestAssociationMap.dev.cc new file mode 100644 index 0000000000000..cf7337eaa56dd --- /dev/null +++ b/DataFormats/TICL/test/alpaka/TestAssociationMap.dev.cc @@ -0,0 +1,280 @@ +#include "DataFormats/Portable/interface/PortableCollection.h" +#include "DataFormats/Portable/interface/PortableHostCollection.h" +#include "DataFormats/TICL/interface/AssociationMap.h" +#include "DataFormats/TICL/interface/FillAssociator.h" + +#include +#include + +#define CATCH_CONFIG_MAIN +#include + +using namespace ALPAKA_ACCELERATOR_NAMESPACE; + +namespace { + template + struct TestKernel { + ALPAKA_FN_ACC void operator()(const Acc1D& acc, ticl::AssociationMapView map, bool* result) const { + for (auto idx : alpaka::uniformElements(acc, map.size() / 2)) { + *result &= (map[0][idx] % 2 != 0); + *result &= (map[1][idx] % 2 == 0); + } + *result &= (map.count(0) == 50); + *result &= (map.count(1) == 50); + *result &= map.contains(0); + *result &= map.contains(1); + *result &= (map.keys() == 2); + } + }; +} // namespace + +TEST_CASE("Construct and fill an association map") { + const auto& devices = cms::alpakatools::devices(); + + SECTION("Test default map") { + for (const auto& device : devices) { + Queue queue(device); + + const auto nkeys = 2u; + const auto nvalues = 100u; + PortableCollection> map(queue, nvalues, nkeys); + auto host_values = cms::alpakatools::make_host_buffer(queue, nvalues); + auto host_associations = cms::alpakatools::make_host_buffer(queue, nvalues); + std::ranges::copy(std::views::iota(0u, nvalues), host_values.data()); + for (auto i : std::views::iota(0u, nvalues)) + host_associations[i] = (i % 2 == 0); + auto device_values = cms::alpakatools::make_device_buffer(queue, nvalues); + auto device_associations = cms::alpakatools::make_device_buffer(queue, nvalues); + alpaka::memcpy(queue, device_values, host_values); + alpaka::memcpy(queue, device_associations, host_associations); + ticl::associator::fill( + queue, map.view(), std::span{device_associations}, std::span{device_values}); + auto offsets = cms::alpakatools::make_host_buffer(queue, map.view().offsets().keys_offsets().size()); + auto values = cms::alpakatools::make_host_buffer(queue, nvalues); + alpaka::memcpy(queue, offsets, cms::alpakatools::make_device_view(queue, map.view().offsets().keys_offsets())); + alpaka::memcpy(queue, values, cms::alpakatools::make_device_view(queue, map.view().content().values())); + + auto device_result = cms::alpakatools::make_device_buffer(queue); + alpaka::memset(queue, device_result, 1); + const auto blocksize = 1024u; + const auto gridsize = cms::alpakatools::divide_up_by(nvalues, blocksize); + auto work_division = cms::alpakatools::make_workdiv(gridsize, blocksize); + alpaka::exec(queue, work_division, TestKernel{}, map.view(), device_result.data()); + auto host_result = cms::alpakatools::make_host_buffer(queue); + alpaka::memcpy(queue, host_result, device_result); + alpaka::wait(queue); + + CHECK(*host_result); + + // check content back on CPU + CHECK(offsets[0] == 0); + CHECK(offsets[1] == 50); + CHECK(offsets[2] == 100); + for (auto i = 0u; i < nvalues; ++i) { + if (i < nvalues / 2) { + CHECK(values[i] % 2 != 0); + } else { + CHECK(values[i] % 2 == 0); + } + } + SECTION("Test partially filled map") { + PortableCollection> partial_map(queue, nvalues, nkeys); + auto partial_map_view = partial_map.view(); + ticl::associator::fill(queue, + partial_map_view, + std::span{device_associations.data(), nvalues / 2}, + std::span{device_values.data(), nvalues / 2}); + auto partial_offsets = + cms::alpakatools::make_host_buffer(queue, partial_map.view().offsets().keys_offsets().size()); + auto partial_values = cms::alpakatools::make_host_buffer(queue, nvalues / 2); + alpaka::memcpy(queue, + partial_offsets, + cms::alpakatools::make_device_view(queue, partial_map.view().offsets().keys_offsets())); + alpaka::memcpy(queue, + partial_values, + cms::alpakatools::make_device_view(queue, partial_map.view().content().values(), nvalues / 2)); + alpaka::wait(queue); + + CHECK(partial_offsets[0] == 0); + CHECK(partial_offsets[1] == nvalues / 4); + CHECK(partial_offsets[2] == nvalues / 2); + for (auto i = 0u; i < nvalues / 2; ++i) { + if (i < nvalues / 4) { + CHECK(partial_values[i] % 2 != 0); + } else { + CHECK(partial_values[i] % 2 == 0); + } + } + } + SECTION("Test deep copy") { + PortableCollection> new_map(queue, nvalues, nkeys); + new_map.deepCopy(queue, map.const_view()); + + alpaka::memset(queue, offsets, 0x0); + alpaka::memset(queue, values, 0x0); + alpaka::memcpy( + queue, offsets, cms::alpakatools::make_device_view(queue, new_map.view().offsets().keys_offsets())); + alpaka::memcpy(queue, values, cms::alpakatools::make_device_view(queue, new_map.view().content().values())); + alpaka::wait(queue); + + CHECK(offsets[0] == 0); + CHECK(offsets[1] == 50); + CHECK(offsets[2] == 100); + for (auto i = 0u; i < nvalues; ++i) { + if (i < nvalues / 2) { + CHECK(values[i] % 2 != 0); + } else { + CHECK(values[i] % 2 == 0); + } + } + } + } + } + SECTION("Test map with specialized templates") { + for (const auto& device : devices) { + Queue queue(device); + + const auto nkeys = 2; + const auto nvalues = 100; + PortableCollection> map(queue, nvalues, nkeys); + auto host_values = cms::alpakatools::make_host_buffer(queue, nvalues); + auto host_associations = cms::alpakatools::make_host_buffer(queue, nvalues); + std::ranges::copy(std::views::iota(0, nvalues), host_values.data()); + for (auto i : std::views::iota(0, nvalues)) + host_associations[i] = (i % 2 == 0); + auto device_values = cms::alpakatools::make_device_buffer(queue, nvalues); + auto device_associations = cms::alpakatools::make_device_buffer(queue, nvalues); + alpaka::memcpy(queue, device_values, host_values); + alpaka::memcpy(queue, device_associations, host_associations); + ticl::associator::fill(queue, + map.view(), + std::span{device_associations.data(), static_cast(nvalues)}, + std::span{device_values.data(), static_cast(nvalues)}); + auto offsets = cms::alpakatools::make_host_buffer(queue, map.view().offsets().keys_offsets().size()); + auto values = cms::alpakatools::make_host_buffer(queue, nvalues); + alpaka::memcpy(queue, offsets, cms::alpakatools::make_device_view(queue, map.view().offsets().keys_offsets())); + alpaka::memcpy(queue, values, cms::alpakatools::make_device_view(queue, map.view().content().values())); + + auto device_result = cms::alpakatools::make_device_buffer(queue); + alpaka::memset(queue, device_result, 1); + const auto blocksize = 1024u; + const auto gridsize = cms::alpakatools::divide_up_by(nvalues, blocksize); + auto work_division = cms::alpakatools::make_workdiv(gridsize, blocksize); + alpaka::exec(queue, work_division, TestKernel{}, map.view(), device_result.data()); + auto host_result = cms::alpakatools::make_host_buffer(queue); + alpaka::memcpy(queue, host_result, device_result); + alpaka::wait(queue); + + CHECK(*host_result); + + // check content back on CPU + CHECK(offsets[0] == 0); + CHECK(offsets[1] == 50); + CHECK(offsets[2] == 100); + for (auto i = 0u; i < nvalues; ++i) { + if (i < nvalues / 2) { + CHECK(values[i] % 2 != 0); + } else { + CHECK(values[i] % 2 == 0); + } + } + SECTION("Test deep copy") { + PortableCollection> new_map(queue, nvalues, nkeys); + new_map.deepCopy(queue, map.const_view()); + + alpaka::memset(queue, offsets, 0x0); + alpaka::memset(queue, values, 0x0); + alpaka::memcpy( + queue, offsets, cms::alpakatools::make_device_view(queue, new_map.view().offsets().keys_offsets())); + alpaka::memcpy(queue, values, cms::alpakatools::make_device_view(queue, new_map.view().content().values())); + alpaka::wait(queue); + + CHECK(offsets[0] == 0); + CHECK(offsets[1] == 50); + CHECK(offsets[2] == 100); + for (auto i = 0u; i < nvalues; ++i) { + if (i < nvalues / 2) { + CHECK(values[i] % 2 != 0); + } else { + CHECK(values[i] % 2 == 0); + } + } + } + } + } +} + +TEST_CASE("Test different numbers of keys") { + const auto& devices = cms::alpakatools::devices(); + + SECTION("Test map with keys close to alignments") { + for (const auto& device : devices) { + Queue queue(device); + + const auto nvalues = 100u; + auto host_values = cms::alpakatools::make_host_buffer(queue, nvalues); + std::ranges::copy(std::views::iota(0u, nvalues), host_values.data()); + + const auto nkeys1 = 30; + const auto nkeys2 = 31; + const auto nkeys3 = 32; + auto host_associations1 = cms::alpakatools::make_host_buffer(queue, nvalues); + auto host_associations2 = cms::alpakatools::make_host_buffer(queue, nvalues); + auto host_associations3 = cms::alpakatools::make_host_buffer(queue, nvalues); + for (auto i : std::views::iota(0u, nvalues)) { + host_associations1[i] = (i < static_cast(nkeys1)) ? i : nkeys1 - 1; + host_associations2[i] = (i < static_cast(nkeys2)) ? i : nkeys2 - 1; + host_associations3[i] = (i < static_cast(nkeys3)) ? i : nkeys3 - 1; + } + auto device_values = cms::alpakatools::make_device_buffer(queue, nvalues); + auto device_associations1 = cms::alpakatools::make_device_buffer(queue, nvalues); + auto device_associations2 = cms::alpakatools::make_device_buffer(queue, nvalues); + auto device_associations3 = cms::alpakatools::make_device_buffer(queue, nvalues); + alpaka::memcpy(queue, device_values, host_values); + alpaka::memcpy(queue, device_associations1, host_associations1); + alpaka::memcpy(queue, device_associations2, host_associations2); + alpaka::memcpy(queue, device_associations3, host_associations3); + + PortableCollection> map1(queue, nvalues, nkeys1); + PortableCollection> map2(queue, nvalues, nkeys2); + PortableCollection> map3(queue, nvalues, nkeys3); + ticl::associator::fill(queue, + map1.view(), + std::span{device_associations1}, + std::span{device_values}); + ticl::associator::fill(queue, + map2.view(), + std::span{device_associations2}, + std::span{device_values}); + ticl::associator::fill(queue, + map3.view(), + std::span{device_associations3}, + std::span{device_values}); + + // every key below `nkeys - 1` gets exactly the single value equal to its own index; + // all the overflow values (i >= nkeys - 1) are collected, in order, under the last key + auto checkMap = [&](auto& map, int nkeys) { + auto offsets = + cms::alpakatools::make_host_buffer(queue, map.view().offsets().keys_offsets().size()); + auto values = cms::alpakatools::make_host_buffer(queue, nvalues); + alpaka::memcpy(queue, offsets, cms::alpakatools::make_device_view(queue, map.view().offsets().keys_offsets())); + alpaka::memcpy(queue, values, cms::alpakatools::make_device_view(queue, map.view().content().values())); + alpaka::wait(queue); + + const auto last_key = static_cast(nkeys - 1); + CHECK(offsets[0] == 0); + for (auto k = 0u; k < last_key; ++k) { + CHECK(offsets[k + 1] - offsets[k] == 1); + CHECK(values[offsets[k]] == k); + } + CHECK(offsets[nkeys] == nvalues); + for (auto i = last_key; i < nvalues; ++i) { + CHECK(values[offsets[last_key] + (i - last_key)] >= last_key); + } + }; + checkMap(map1, nkeys1); + checkMap(map2, nkeys2); + checkMap(map3, nkeys3); + } + } +}