Commit a8994fe9 authored by Roel Aaij's avatar Roel Aaij
Browse files

Create Consumer for all VP non-event-data.

parent 95a18454
......@@ -71,20 +71,16 @@ struct VeloRawBank {
* @brief Velo geometry description typecast.
*/
struct VeloGeometry {
size_t size;
float* local_x;
float* x_pitch;
float* ltg;
size_t n_trans;
float module_zs[Velo::Constants::n_modules];
float local_x[Velo::Constants::number_of_sensor_columns];
float x_pitch[Velo::Constants::number_of_sensor_columns];
float ltg[12 * Velo::Constants::n_sensors];
/**
* @brief Typecast from std::vector.
*/
VeloGeometry(const std::vector<char>& geometry);
/**
* @brief Just typecast, no size check.
*/
__device__ __host__ VeloGeometry(const char* geometry);
};
__device__ __host__ uint32_t get_channel_id(uint sensor, uint chip, uint col, uint row);
......
#include <cassert>
#include "ClusteringDefinitions.cuh"
__device__ __host__ VeloRawEvent::VeloRawEvent(const char* event)
......@@ -20,38 +21,41 @@ __device__ __host__ VeloRawBank::VeloRawBank(const char* raw_bank)
sp_word = (uint32_t*) p;
}
VeloGeometry::VeloGeometry(const std::vector<char>& geometry)
VeloGeometry::VeloGeometry(std::vector<char> const& geometry)
{
const char* p = geometry.data();
char const* p = geometry.data();
local_x = (float*) p;
p += sizeof(float) * Velo::Constants::number_of_sensor_columns;
x_pitch = (float*) p;
p += sizeof(float) * Velo::Constants::number_of_sensor_columns;
ltg = (float*) p;
p += sizeof(float) * 16 * Velo::Constants::n_sensors;
auto copy_array = [this, &p] (const size_t N, float* d) {
const size_t n = ((size_t*)p)[0];
if (n != N) {
error_cout << n << " != " << N << std::endl;
}
p += sizeof(size_t);
memcpy(d, p, sizeof(float) * n);
p += sizeof(float) * n;
};
size = p - geometry.data();
copy_array(Velo::Constants::n_modules, module_zs);
copy_array(Velo::Constants::number_of_sensor_columns, local_x);
copy_array(Velo::Constants::number_of_sensor_columns, x_pitch);
size_t n_ltg = ((size_t*)p)[0];
assert(n_ltg == Velo::Constants::n_sensors);
p += sizeof(size_t);
n_trans = ((size_t*)p)[0];
assert(n_trans == 12);
p += sizeof(size_t);
for (size_t i = 0; i < n_ltg; ++i) {
memcpy(ltg + n_trans * i, p, n_trans * sizeof(float));
p += sizeof(float) * n_trans;
}
const size_t size = p - geometry.data();
if (size != geometry.size()) {
error_cout << "Size mismatch for geometry" << std::endl;
}
}
__device__ __host__ VeloGeometry::VeloGeometry(const char* geometry)
{
const char* p = geometry;
local_x = (float*) p;
p += sizeof(float) * Velo::Constants::number_of_sensor_columns;
x_pitch = (float*) p;
p += sizeof(float) * Velo::Constants::number_of_sensor_columns;
ltg = (float*) p;
p += sizeof(float) * 16 * Velo::Constants::n_sensors;
size = p - geometry;
}
__device__ __host__ uint32_t get_channel_id(unsigned int sensor, unsigned int chip, unsigned int col, unsigned int row)
{
return (sensor << LHCb::VPChannelID::sensorBits) | (chip << LHCb::VPChannelID::chipBits) |
......
......@@ -17,7 +17,7 @@ __global__ void masked_velo_clustering(
uint32_t* dev_velo_cluster_container,
const uint* dev_event_list,
uint* dev_event_order,
char* dev_velo_geometry,
const VeloGeometry* dev_velo_geometry,
uint8_t* dev_velo_sp_patterns,
float* dev_velo_sp_fx,
float* dev_velo_sp_fy);
......
......@@ -24,7 +24,7 @@ __global__ void masked_velo_clustering(
uint32_t* dev_velo_cluster_container,
const uint* dev_event_list,
uint* dev_event_order,
char* dev_velo_geometry,
const VeloGeometry* dev_velo_geometry,
uint8_t* dev_velo_sp_patterns,
float* dev_velo_sp_fx,
float* dev_velo_sp_fy)
......@@ -45,7 +45,7 @@ __global__ void masked_velo_clustering(
float* float_velo_cluster_container = (float*) dev_velo_cluster_container;
// Load Velo geometry (assume it is the same for all events)
const VeloGeometry g(dev_velo_geometry);
const VeloGeometry& g = *dev_velo_geometry;
// Read raw event
const auto raw_event = VeloRawEvent(raw_input);
......@@ -57,7 +57,7 @@ __global__ void masked_velo_clustering(
// Read raw bank
const auto raw_bank = VeloRawBank(raw_event.payload + raw_event.raw_bank_offset[raw_bank_number]);
const float* ltg = g.ltg + 16 * raw_bank.sensor_index;
const float* ltg = g.ltg + g.n_trans * raw_bank.sensor_index;
for (int sp_index = 0; sp_index < raw_bank.sp_count; ++sp_index) {
// Decode sp
......@@ -147,7 +147,7 @@ __global__ void masked_velo_clustering(
assert(raw_bank_number < Velo::Constants::n_sensors);
const auto raw_bank = VeloRawBank(raw_event.payload + raw_event.raw_bank_offset[raw_bank_number]);
const float* ltg = g.ltg + 16 * raw_bank.sensor_index;
const float* ltg = g.ltg + g.n_trans * raw_bank.sensor_index;
const uint32_t sp_word = raw_bank.sp_word[sp_index];
const uint32_t sp_addr = (sp_word & 0x007FFF00U) >> 8;
// Note: In the code below, row and col are int32_t (not unsigned)
......
......@@ -2,6 +2,7 @@
#include <cstdint>
#include <cfloat>
#include "ClusteringDefinitions.cuh"
#include "VeloEventModel.cuh"
#include "FillCandidates.cuh"
#include "ProcessModules.cuh"
......@@ -24,7 +25,7 @@ __global__ void search_by_triplet(
short* dev_h0_candidates,
short* dev_h2_candidates,
unsigned short* dev_rel_indices,
const float* dev_velo_module_zs);
const VeloGeometry* dev_velo_geometry);
ALGORITHM(
search_by_triplet,
......
......@@ -42,7 +42,7 @@ __global__ void search_by_triplet(
short* dev_h0_candidates,
short* dev_h2_candidates,
unsigned short* dev_rel_indices,
const float* dev_velo_module_zs)
const VeloGeometry* dev_velo_geometry)
{
/* Data initialization */
// Each event is treated with two blocks, one for each side.
......@@ -88,6 +88,6 @@ __global__ void search_by_triplet(
number_of_hits,
h1_rel_indices,
hit_offset,
dev_velo_module_zs,
dev_velo_geometry->module_zs,
dev_atomics_velo);
}
......@@ -27,6 +27,17 @@ namespace Consumers {
std::reference_wrapper<gsl::span<char>> m_dev_geometry;
};
struct VPGeometry final : public Allen::NonEventData::Consumer {
public:
VPGeometry(Constants& constants);
void consume(std::vector<char> const& data) override;
private:
void initialize(const std::vector<char>& data);
std::reference_wrapper<Constants> m_constants;
};
struct UTGeometry final : public Allen::NonEventData::Consumer {
public:
UTGeometry(Constants& constants);
......@@ -62,24 +73,22 @@ namespace Consumers {
struct Beamline final : public Allen::NonEventData::Consumer {
public:
Beamline(float*&);
Beamline(gsl::span<float>&);
void consume(std::vector<char> const& data) override;
private:
std::reference_wrapper<float*> m_dev_beamline;
const size_t m_size = 2 * sizeof(float);
std::reference_wrapper<gsl::span<float>> m_dev_beamline;
};
struct MagneticField final : public Allen::NonEventData::Consumer {
public:
MagneticField(float*&);
MagneticField(gsl::span<float>&);
void consume(std::vector<char> const& data) override;
private:
std::reference_wrapper<float*> m_dev_magnet_polarity;
const size_t m_size = sizeof(float);
std::reference_wrapper<gsl::span<float>> m_dev_magnet_polarity;
};
struct MuonGeometry final : public Allen::NonEventData::Consumer {
......@@ -103,9 +112,9 @@ namespace Consumers {
static constexpr size_t n_data_blocks = 27;
MuonLookupTables(
std::vector<char>& host_muon_tables_raw,
char*& dev_muon_tables_raw,
Muon::MuonTables*& muon_tables);
std::vector<char>& host_muon_tables_raw,
char*& dev_muon_tables_raw,
Muon::MuonTables*& muon_tables);
void consume(std::vector<char> const& data) override;
......
......@@ -9,17 +9,20 @@ namespace {
using std::to_string;
} // namespace
Consumers::Beamline::Beamline(float*& dev_beamline) : m_dev_beamline {dev_beamline} {}
Consumers::Beamline::Beamline(gsl::span<float>& dev_beamline) : m_dev_beamline {dev_beamline} {}
void Consumers::Beamline::consume(std::vector<char> const& data)
{
if (data.size() != m_size) {
throw StrException {string {"sizes don't match: "} + to_string(m_size) + " " + to_string(data.size())};
}
if (!m_dev_beamline.get()) {
if (m_dev_beamline.get().empty()) {
// Allocate space
char* p = nullptr;
cudaCheck(cudaMalloc((void**) &m_dev_beamline.get(), data.size()));
float* p = nullptr;
cudaCheck(cudaMalloc((void**) &p, data.size()));
m_dev_beamline.get() = {p, data.size() / sizeof(float)};
}
else if (data.size() != m_dev_beamline.get().size()) {
throw StrException {string {"sizes don't match: "} + to_string(m_dev_beamline.get().size()) + " " +
to_string(data.size())};
}
cudaCheck(cudaMemcpy(m_dev_beamline.get(), data.data(), data.size(), cudaMemcpyHostToDevice));
cudaCheck(cudaMemcpy(m_dev_beamline.get().data(), data.data(), data.size(), cudaMemcpyHostToDevice));
}
......@@ -9,17 +9,22 @@ namespace {
using std::to_string;
} // namespace
Consumers::MagneticField::MagneticField(float*& dev_magnet_polarity) : m_dev_magnet_polarity {dev_magnet_polarity} {}
Consumers::MagneticField::MagneticField(gsl::span<float>& dev_magnet_polarity) :
m_dev_magnet_polarity {dev_magnet_polarity}
{}
void Consumers::MagneticField::consume(std::vector<char> const& data)
{
if (data.size() != m_size) {
throw StrException {string {"sizes don't match: "} + to_string(m_size) + " " + to_string(data.size())};
}
if (!m_dev_magnet_polarity.get()) {
if (m_dev_magnet_polarity.get().empty()) {
// Allocate space
char* p = nullptr;
cudaCheck(cudaMalloc((void**) &m_dev_magnet_polarity.get(), data.size()));
float* p = nullptr;
cudaCheck(cudaMalloc((void**) &p, data.size()));
m_dev_magnet_polarity.get() = {p, data.size() / sizeof(float)};
}
else if (data.size() != m_dev_magnet_polarity.get().size()) {
throw StrException {string {"sizes don't match: "} + to_string(m_dev_magnet_polarity.get().size()) + " " +
to_string(data.size())};
}
cudaCheck(cudaMemcpy(m_dev_magnet_polarity, data.data(), data.size(), cudaMemcpyHostToDevice));
cudaCheck(cudaMemcpy(m_dev_magnet_polarity.get().data(), data.data(), data.size(), cudaMemcpyHostToDevice));
}
#include <string>
#include <cuda_runtime.h>
#include <Common.h>
#include <Consumers.h>
namespace {
using std::string;
using std::to_string;
using std::vector;
} // namespace
Consumers::VPGeometry::VPGeometry(Constants& constants) : m_constants {constants} {}
void Consumers::VPGeometry::initialize(vector<char> const& data)
{
auto alloc_and_copy = [](auto const& host_numbers, auto& device_numbers) {
using value_type = typename std::remove_reference_t<decltype(host_numbers)>::value_type;
value_type* p = nullptr;
cudaCheck(cudaMalloc((void**) &p, host_numbers.size() * sizeof(value_type)));
device_numbers = gsl::span {p, host_numbers.size()};
cudaCheck(cudaMemcpy(
device_numbers.data(), host_numbers.data(), host_numbers.size() * sizeof(value_type), cudaMemcpyHostToDevice));
};
// Velo clustering candidate ks
std::array<uint8_t, VeloClustering::lookup_table_size> host_candidate_ks = {0, 0, 1, 4, 4, 5, 5, 5, 5};
alloc_and_copy(host_candidate_ks, m_constants.get().dev_velo_candidate_ks);
// Velo clustering patterns
// Fetch patterns and populate in GPU
vector<uint8_t> sp_patterns(256, 0);
vector<uint8_t> sp_sizes(256, 0);
vector<float> sp_fx(512, 0);
vector<float> sp_fy(512, 0);
cache_sp_patterns(sp_patterns, sp_sizes, sp_fx, sp_fy);
alloc_and_copy(sp_patterns, m_constants.get().dev_velo_sp_patterns);
alloc_and_copy(sp_fx, m_constants.get().dev_velo_sp_fx);
alloc_and_copy(sp_fy, m_constants.get().dev_velo_sp_fy);
cudaCheck(cudaMalloc((void**) &m_constants.get().dev_velo_geometry, sizeof(VeloGeometry)));
}
void Consumers::VPGeometry::consume(vector<char> const& data)
{
auto& dev_velo_geometry = m_constants.get().dev_velo_geometry;
if (dev_velo_geometry == nullptr) {
initialize(data);
}
else if (sizeof(VeloGeometry) != data.size()) {
throw StrException {string {"sizes don't match: "} + to_string(sizeof(VeloGeometry)) + " " +
to_string(data.size())};
}
VeloGeometry host_velo_geometry {data};
cudaCheck(cudaMemcpy(dev_velo_geometry, &host_velo_geometry, sizeof(VeloGeometry), cudaMemcpyHostToDevice));
}
......@@ -49,7 +49,7 @@ void register_consumers(Allen::NonEventData::IUpdater* updater, Constants& const
tuple {Allen::NonEventData::MagneticField {},
std::make_unique<Consumers::MagneticField>(constants.dev_magnet_polarity)},
tuple {Allen::NonEventData::Beamline {}, std::make_unique<Consumers::Beamline>(constants.dev_beamline)},
tuple {Allen::NonEventData::VeloGeometry {}, std::make_unique<Consumers::RawGeometry>(constants.dev_velo_geometry)},
tuple {Allen::NonEventData::VeloGeometry{}, std::make_unique<Consumers::VPGeometry>(constants)},
tuple {Allen::NonEventData::MuonGeometry {},
std::make_unique<Consumers::MuonGeometry>(
constants.host_muon_geometry_raw, constants.dev_muon_geometry_raw, constants.dev_muon_geometry)},
......
template <typename Derived>
class GPUAlgorithm{};
......@@ -32,15 +32,11 @@
*/
struct Constants {
std::array<uint8_t, VeloClustering::lookup_table_size> host_candidate_ks;
std::array<float, 9> host_inv_clus_res;
float* dev_velo_module_zs = nullptr;
uint8_t* dev_velo_candidate_ks = nullptr;
uint8_t* dev_velo_sp_patterns = nullptr;
float* dev_velo_sp_fx = nullptr;
float* dev_velo_sp_fy = nullptr;
char* dev_velo_geometry = nullptr;
gsl::span<uint8_t> dev_velo_candidate_ks;
gsl::span<uint8_t> dev_velo_sp_patterns;
gsl::span<float> dev_velo_sp_fx;
gsl::span<float> dev_velo_sp_fy;
VeloGeometry* dev_velo_geometry = nullptr;
std::vector<char> host_ut_geometry;
std::array<uint, UT::Constants::n_layers * UT::Constants::n_regions_in_layer + 1> host_ut_region_offsets;
......@@ -58,6 +54,7 @@ struct Constants {
gsl::span<char> dev_ut_boards;
UTMagnetTool* dev_ut_magnet_tool = nullptr;
std::array<float, 9> host_inv_clus_res;
SciFi::Tracking::TMVA* dev_scifi_tmva1 = nullptr;
SciFi::Tracking::TMVA* dev_scifi_tmva2 = nullptr;
SciFi::Tracking::Arrays* dev_scifi_constArrays;
......@@ -68,10 +65,10 @@ struct Constants {
std::vector<char> host_scifi_geometry;
// Beam location
float* dev_beamline = nullptr;
gsl::span<float> dev_beamline;
// Magnet polarity
float* dev_magnet_polarity = nullptr;
gsl::span<float> dev_magnet_polarity;
// Looking forward
LookingForward::Constants host_looking_forward_constants;
......
......@@ -2,11 +2,6 @@
void Constants::reserve_constants()
{
cudaCheck(cudaMalloc((void**) &dev_velo_module_zs, Velo::Constants::n_modules * sizeof(float)));
cudaCheck(cudaMalloc((void**) &dev_velo_candidate_ks, 9 * sizeof(uint8_t)));
cudaCheck(cudaMalloc((void**) &dev_velo_sp_patterns, 256 * sizeof(uint8_t)));
cudaCheck(cudaMalloc((void**) &dev_velo_sp_fx, 512 * sizeof(float)));
cudaCheck(cudaMalloc((void**) &dev_velo_sp_fy, 512 * sizeof(float)));
cudaCheck(cudaMalloc((void**) &dev_scifi_tmva1, sizeof(SciFi::Tracking::TMVA)));
cudaCheck(cudaMalloc((void**) &dev_scifi_tmva2, sizeof(SciFi::Tracking::TMVA)));
cudaCheck(cudaMalloc((void**) &dev_scifi_constArrays, sizeof(SciFi::Tracking::Arrays)));
......@@ -15,8 +10,6 @@ void Constants::reserve_constants()
cudaCheck(cudaMalloc((void**) &dev_looking_forward_constants, sizeof(LookingForward::Constants)));
cudaCheck(cudaMalloc((void**) &dev_muon_foi, sizeof(Muon::Constants::FieldOfInterest)));
cudaCheck(cudaMalloc((void**) &dev_muon_momentum_cuts, 3 * sizeof(float)));
cudaCheck(cudaMalloc((void**) &dev_magnet_polarity, sizeof(float)));
cudaCheck(cudaMalloc((void**) &dev_beamline, 2 * sizeof(float)));
}
void Constants::initialize_constants(
......@@ -25,33 +18,13 @@ void Constants::initialize_constants(
{
// Velo module constants
const std::array<float, Velo::Constants::n_modules> velo_module_zs = {
-287.5, -275, -262.5, -250, -237.5, -225, -212.5, -200, -137.5, -125, -62.5, -50, -37.5,
-25, -12.5, 0, 12.5, 25, 37.5, 50, 62.5, 75, 87.5, 100, 112.5, 125,
137.5, 150, 162.5, 175, 187.5, 200, 212.5, 225, 237.5, 250, 262.5, 275, 312.5,
325, 387.5, 400, 487.5, 500, 587.5, 600, 637.5, 650, 687.5, 700, 737.5, 750};
cudaCheck(cudaMemcpy(
dev_velo_module_zs, velo_module_zs.data(), velo_module_zs.size() * sizeof(float), cudaMemcpyHostToDevice));
// Velo clustering candidate ks
host_candidate_ks = {0, 0, 1, 4, 4, 5, 5, 5, 5};
cudaCheck(cudaMemcpy(
dev_velo_candidate_ks,
host_candidate_ks.data(),
host_candidate_ks.size() * sizeof(uint8_t),
cudaMemcpyHostToDevice));
// Velo clustering patterns
// Fetch patterns and populate in GPU
std::vector<uint8_t> sp_patterns(256, 0);
std::vector<uint8_t> sp_sizes(256, 0);
std::vector<float> sp_fx(512, 0);
std::vector<float> sp_fy(512, 0);
cache_sp_patterns(sp_patterns, sp_sizes, sp_fx, sp_fy);
cudaCheck(cudaMemcpy(dev_velo_sp_patterns, sp_patterns.data(), sp_patterns.size(), cudaMemcpyHostToDevice));
cudaCheck(cudaMemcpy(dev_velo_sp_fx, sp_fx.data(), sp_fx.size() * sizeof(float), cudaMemcpyHostToDevice));
cudaCheck(cudaMemcpy(dev_velo_sp_fy, sp_fy.data(), sp_fy.size() * sizeof(float), cudaMemcpyHostToDevice));
// const std::array<float, Velo::Constants::n_modules> velo_module_zs = {
// -287.5, -275, -262.5, -250, -237.5, -225, -212.5, -200, -137.5, -125, -62.5, -50, -37.5,
// -25, -12.5, 0, 12.5, 25, 37.5, 50, 62.5, 75, 87.5, 100, 112.5, 125,
// 137.5, 150, 162.5, 175, 187.5, 200, 212.5, 225, 237.5, 250, 262.5, 275, 312.5,
// 325, 387.5, 400, 487.5, 500, 587.5, 600, 637.5, 650, 687.5, 700, 737.5, 750};
// cudaCheck(cudaMemcpy(
// dev_velo_module_zs, velo_module_zs.data(), velo_module_zs.size() * sizeof(float), cudaMemcpyHostToDevice));
// SciFi constants
SciFi::Tracking::TMVA host_tmva1;
......
......@@ -43,7 +43,7 @@ void SequenceVisitor::visit<lf_fit_t>(
constants.dev_inv_clus_res,
constants.dev_scifi_constArrays,
constants.dev_looking_forward_constants,
constants.dev_magnet_polarity,
constants.dev_magnet_polarity.data(),
arguments.offset<dev_scifi_lf_track_params>());
state.invoke();
......
......@@ -53,7 +53,7 @@ void SequenceVisitor::visit<lf_quality_filter_t>(
constants.dev_scifi_tmva2,
constants.dev_scifi_constArrays,
constants.dev_looking_forward_constants,
constants.dev_magnet_polarity,
constants.dev_magnet_polarity.data(),
arguments.offset<dev_atomics_scifi>(),
arguments.offset<dev_scifi_selected_track_indices>(),
arguments.offset<dev_scifi_tracks>());
......
......@@ -44,7 +44,7 @@ void SequenceVisitor::visit<lf_search_initial_windows_t>(
constants.dev_scifi_geometry,
constants.dev_inv_clus_res,
constants.dev_scifi_constArrays,
constants.dev_magnet_polarity,
constants.dev_magnet_polarity.data(),
constants.dev_looking_forward_constants,
arguments.offset<dev_scifi_lf_initial_windows>(),
arguments.offset<dev_ut_states>());
......
......@@ -34,7 +34,7 @@ void SequenceVisitor::visit<compass_ut_t>(
arguments.offset<dev_velo_track_hits>(),
arguments.offset<dev_velo_states>(),
constants.dev_ut_magnet_tool,
constants.dev_magnet_polarity,
constants.dev_magnet_polarity.data(),
constants.dev_ut_dxDy.data(),
arguments.offset<dev_ut_active_tracks>(),
constants.dev_unique_x_sector_layer_offsets.data(),
......
......@@ -29,7 +29,7 @@ void SequenceVisitor::visit<pv_beamline_histo_t>(
arguments.offset<dev_velo_track_hit_number>(),
arguments.offset<dev_pvtracks>(),
arguments.offset<dev_zhisto>(),
constants.dev_beamline);
constants.dev_beamline.data());
state.invoke();
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment