Skip to content
Snippets Groups Projects

Compare revisions

Changes are shown as if the source revision was being merged into the target revision. Learn more about comparing revisions.

Source

Select target project
No results found

Target

Select target project
  • pvfinder/inference-engine
  • rmatev/Allen
  • suali/Allen
  • mstahl/Allen
  • roneil/Allen
  • graemes/Allen
  • cburr/Allen
  • jonrob/Allen
  • bjashal/Allen-HIP
  • dcampora/MiniAllen
  • brij/Allen
  • raaij/cuda_hlt
  • bsm-fleet/cuda_hlt
  • abrearod/cuda_hlt
  • aalvesju/cuda_hlt
  • lhcb/Allen
16 results
Show changes
Commits on Source (33)
Showing
with 783 additions and 451 deletions
......@@ -58,6 +58,7 @@ parser.add_argument(
default=os.path.join(allen_dir, "input", "detector_configuration"))
parser.add_argument("-n", dest="n_events", default=0)
parser.add_argument("-t", dest="threads", default=1)
parser.add_argument("--params", dest="params", default="")
parser.add_argument("-r", dest="repetitions", default=1)
parser.add_argument("-m", dest="reserve", default=1024)
parser.add_argument("-v", dest="verbosity", default=3)
......@@ -270,8 +271,9 @@ zmqSvc = gaudi.service("ZeroMQSvc", interface=gbl.IZeroMQSvc)
# options map
options = gbl.std.map("std::string", "std::string")()
for flag, value in [("g", args.det_folder),
("params", os.getenv("PARAMFILESROOT")),
params = args.params if args.params != "" else os.getenv("PARAMFILESROOT")
for flag, value in [("g", args.det_folder), ("params", params),
("n", args.n_events), ("t", args.threads),
("r", args.repetitions), ("output-file", args.output_file),
("output-batch-size", args.output_batch_size),
......
......@@ -24,7 +24,12 @@ file(MAKE_DIRECTORY ${ALLEN_GENERATED_INCLUDE_FILES_DIR})
file(MAKE_DIRECTORY ${ALLEN_ALGORITHMS_DIR})
# We will invoke the parser a few times, set its required environment in a variable
set(PARSER_ENV PYTHONPATH=$ENV{PYTHONPATH}:${PROJECT_SOURCE_DIR}/scripts LD_LIBRARY_PATH=${LIBCLANG_LIBDIR}:$ENV{LD_LIBRARY_PATH})
# Add the scripts folder only if we are invoking with a CMAKE_TOOLCHAIN_FILE
if(CMAKE_TOOLCHAIN_FILE)
set(PARSER_ENV PYTHONPATH=$ENV{PYTHONPATH}:${PROJECT_SOURCE_DIR}/scripts LD_LIBRARY_PATH=${LIBCLANG_LIBDIR}:$ENV{LD_LIBRARY_PATH})
else()
set(PARSER_ENV LD_LIBRARY_PATH=${LIBCLANG_LIBDIR}:$ENV{LD_LIBRARY_PATH})
endif()
# Parse Allen algorithms
# Parsing should depend on ALL algorithm headers (which include the Parameters section)
......
......@@ -8,6 +8,7 @@ from AllenCore.algorithms import (velo_lumi_counters_t, pv_lumi_counters_t,
calo_lumi_counters_t, plume_lumi_counters_t,
calc_lumi_sum_size_t, make_lumi_summary_t)
from AllenCore.algorithms import muon_calculate_srq_size_t
from AllenCore.configuration_options import allen_register_keys
from AllenConf.odin import decode_odin
from AllenConf.utils import initialize_number_of_events, make_dummy
from AllenCore.generator import make_algorithm
......@@ -15,13 +16,14 @@ from AllenCore.generator import make_algorithm
from AllenConf.persistency import make_gather_selections
from PyConf.tonic import configurable
from PyConf.filecontent_metadata import _get_hash_for_text
from PyConf.filecontent_metadata import register_encoding_dictionary
from AllenConf.velo_reconstruction import decode_velo, make_velo_tracks
from AllenConf.scifi_reconstruction import decode_scifi
from AllenConf.muon_reconstruction import decode_muon
from AllenConf.primary_vertex_reconstruction import make_pvs
from AllenConf.calo_reconstruction import decode_calo
from AllenConf.lumi_schema_generator import LumiSchemaGenerator
from AllenConf.plume_reconstruction import decode_plume
......@@ -40,169 +42,11 @@ def get_lumi_info(lumiInfos, name):
return dummy.dev_lumi_dummy_t
def lumi_summary_maker(lumiInfos, prefix_sum_lumi_size):
def lumi_summary_maker(lumiInfos, prefix_sum_lumi_size, key, lumi_sum_length,
schema):
number_of_events = initialize_number_of_events()
odin = decode_odin()
# TODO generate key here, but it does not save the table anywhere
# so it is not actual usable
table = json.dumps({
"version":
0,
"size":
76,
"counters": [{
"name": "encodingKey",
"offset": 0,
"size": 32
}, {
"name": "T0Low",
"offset": 32,
"size": 32
}, {
"name": "T0High",
"offset": 64,
"size": 32
}, {
"name": "BCIDLow",
"offset": 96,
"size": 32
}, {
"name": "PlumeLumiOverthrLow",
"offset": 128,
"size": 22
}, {
"name": "MuonHitsM3R2",
"offset": 150,
"size": 10
}, {
"name": "PlumeLumiOverthrHigh",
"offset": 160,
"size": 22
}, {
"name": "MuonHitsM4R1",
"offset": 182,
"size": 10
}, {
"name": "SciFiClustersS3M45",
"offset": 192,
"size": 16
}, {
"name": "SciFiClusters",
"offset": 208,
"size": 16
}, {
"name": "SciFiClustersS2M123",
"offset": 224,
"size": 16
}, {
"name": "SciFiClustersS3M123",
"offset": 240,
"size": 16
}, {
"name": "ECalET",
"offset": 256,
"size": 16
}, {
"name": "ECalEInnerTop",
"offset": 272,
"size": 16
}, {
"name": "ECalEMiddleTop",
"offset": 288,
"size": 16
}, {
"name": "ECalEOuterTop",
"offset": 304,
"size": 16
}, {
"name": "ECalEInnerBottom",
"offset": 320,
"size": 16
}, {
"name": "ECalEMiddleBottom",
"offset": 336,
"size": 16
}, {
"name": "ECalEOuterBottom",
"offset": 352,
"size": 16
}, {
"name": "MuonHitsM2R1",
"offset": 368,
"size": 16
}, {
"name": "MuonHitsM2R2",
"offset": 384,
"size": 16
}, {
"name": "MuonHitsM2R3",
"offset": 400,
"size": 16
}, {
"name": "VeloTracks",
"offset": 416,
"size": 15
}, {
"name": "BCIDHigh",
"offset": 431,
"size": 14
}, {
"name": "BXType",
"offset": 445,
"size": 2
}, {
"name": "GEC",
"offset": 447,
"size": 1
}, {
"name": "SciFiClustersS1M45",
"offset": 448,
"size": 13
}, {
"name": "SciFiClustersS2M45",
"offset": 461,
"size": 13
}, {
"name": "VeloVertices",
"offset": 474,
"size": 6
}, {
"name": "PlumeAvgLumiADC",
"offset": 480,
"size": 12
}, {
"name": "MuonHitsM2R4",
"offset": 492,
"size": 11
}, {
"name": "MuonHitsM3R1",
"offset": 512,
"size": 11
}, {
"name": "MuonHitsM3R3",
"offset": 523,
"size": 11
}, {
"name": "MuonHitsM4R4",
"offset": 534,
"size": 10
}, {
"name": "MuonHitsM3R4",
"offset": 544,
"size": 11
}, {
"name": "MuonHitsM4R2",
"offset": 555,
"size": 11
}, {
"name": "MuonHitsM4R3",
"offset": 576,
"size": 11
}]
})
key = int(_get_hash_for_text(table)[:8], 16)
return make_algorithm(
make_lumi_summary_t,
name="make_lumi_summary",
......@@ -217,7 +61,9 @@ def lumi_summary_maker(lumiInfos, prefix_sum_lumi_size):
dev_scifi_info_t=get_lumi_info(lumiInfos, "scifi"),
dev_muon_info_t=get_lumi_info(lumiInfos, "muon"),
dev_calo_info_t=get_lumi_info(lumiInfos, "calo"),
dev_plume_info_t=get_lumi_info(lumiInfos, "plume"))
dev_plume_info_t=get_lumi_info(lumiInfos, "plume"),
lumi_sum_length=lumi_sum_length,
lumi_counter_schema=schema)
def lumi_reconstruction(gather_selections,
......@@ -242,13 +88,50 @@ def lumi_reconstruction(gather_selections,
pvs = make_pvs(velo_tracks)
decoded_muon = decode_muon(empty_banks=not with_muon)
counterSpecs = [("T0Low", 0xffffffff), ("T0High", 0xffffffff),
("BCIDLow", 0xffffffff), ("BCIDHigh", 0x3fff),
("BXType", 3), ("GEC", 1), ("VeloTracks", 1913),
("VeloVertices", 33), ("SciFiClustersS1M45", 765),
("SciFiClustersS2M45", 805), ("SciFiClustersS3M45", 1405),
("SciFiClusters", 7650), ("SciFiClustersS2M123", 7590),
("SciFiClustersS3M123", 7890), ("ECalET", 1072742),
("ECalEInnerTop", 3797317), ("ECalEMiddleTop", 1478032),
("ECalEOuterTop", 1192952), ("ECalEInnerBottom", 4026243),
("ECalEMiddleBottom", 1492195),
("ECalEOuterBottom", 1384124), ("MuonHitsM2R1", 696),
("MuonHitsM2R2", 593), ("MuonHitsM2R3", 263),
("MuonHitsM2R4", 200), ("MuonHitsM3R1", 478),
("MuonHitsM3R2", 212), ("MuonHitsM3R3", 161),
("MuonHitsM3R4", 102), ("MuonHitsM4R1", 134),
("MuonHitsM4R2", 108), ("MuonHitsM4R3", 409),
("MuonHitsM4R4", 227), ("PlumeAvgLumiADC", 0xfff),
("PlumeLumiOverthrLow", 0x3fffff),
("PlumeLumiOverthrHigh", 0x3fffff)]
l = LumiSchemaGenerator(counterSpecs)
l.process()
table = l.getJSON()
if allen_register_keys():
key = int(
register_encoding_dictionary(
"counters", table, directory="luminosity_counters"), 16)
else:
key = 0
lumi_sum_length = table[
"size"] / 4 #algorithms expect length in words not bytes
schema_for_algorithms = {
counter["name"]: (counter["offset"], counter["size"])
for counter in table["counters"]
}
calc_lumi_sum_size = make_algorithm(
calc_lumi_sum_size_t,
name="calc_lumi_sum_size",
host_number_of_events_t=number_of_events["host_number_of_events"],
dev_selections_t=gather_selections.dev_selections_t,
dev_selections_offsets_t=gather_selections.dev_selections_offsets_t,
line_index=lumiLine_index)
line_index=lumiLine_index,
lumi_sum_length=lumi_sum_length)
prefix_sum_lumi_size = make_algorithm(
host_prefix_sum_t,
......@@ -267,7 +150,9 @@ def lumi_reconstruction(gather_selections,
dev_output_buffer_t,
dev_velo_tracks_view_t=velo_tracks["dev_velo_tracks_view"],
dev_offsets_all_velo_tracks_t=velo_tracks[
"dev_offsets_all_velo_tracks"])
"dev_offsets_all_velo_tracks"],
lumi_sum_length=lumi_sum_length,
lumi_counter_schema=schema_for_algorithms)
lumiInfos["pv"] = make_algorithm(
pv_lumi_counters_t,
......@@ -278,7 +163,9 @@ def lumi_reconstruction(gather_selections,
dev_lumi_summary_offsets_t=prefix_sum_lumi_size.
dev_output_buffer_t,
dev_multi_final_vertices_t=pvs["dev_multi_final_vertices"],
dev_number_of_pvs_t=pvs["dev_number_of_multi_final_vertices"])
dev_number_of_pvs_t=pvs["dev_number_of_multi_final_vertices"],
lumi_sum_length=lumi_sum_length,
lumi_counter_schema=schema_for_algorithms)
if with_SciFi:
lumiInfos["scifi"] = make_algorithm(
......@@ -290,7 +177,9 @@ def lumi_reconstruction(gather_selections,
dev_lumi_summary_offsets_t=prefix_sum_lumi_size.
dev_output_buffer_t,
dev_scifi_hit_offsets_t=decoded_scifi["dev_scifi_hit_offsets"],
dev_scifi_hits_t=decoded_scifi["dev_scifi_hits"])
dev_scifi_hits_t=decoded_scifi["dev_scifi_hits"],
lumi_sum_length=lumi_sum_length,
lumi_counter_schema=schema_for_algorithms)
if with_muon:
lumiInfos["muon"] = make_algorithm(
......@@ -302,7 +191,9 @@ def lumi_reconstruction(gather_selections,
dev_lumi_summary_offsets_t=prefix_sum_lumi_size.
dev_output_buffer_t,
dev_storage_station_region_quarter_offsets_t=decoded_muon[
"dev_storage_station_region_quarter_offsets"])
"dev_storage_station_region_quarter_offsets"],
lumi_sum_length=lumi_sum_length,
lumi_counter_schema=schema_for_algorithms)
if with_calo:
lumiInfos["calo"] = make_algorithm(
......@@ -314,7 +205,9 @@ def lumi_reconstruction(gather_selections,
dev_lumi_summary_offsets_t=prefix_sum_lumi_size.
dev_output_buffer_t,
dev_ecal_digits_t=decoded_calo["dev_ecal_digits"],
dev_ecal_digits_offsets_t=decoded_calo["dev_ecal_digits_offsets"])
dev_ecal_digits_offsets_t=decoded_calo["dev_ecal_digits_offsets"],
lumi_sum_length=lumi_sum_length,
lumi_counter_schema=schema_for_algorithms)
if with_plume:
decoded_plume = decode_plume()
......@@ -328,7 +221,9 @@ def lumi_reconstruction(gather_selections,
dev_output_buffer_t,
dev_plume_t=decoded_plume["dev_plume"])
make_lumi_summary = lumi_summary_maker(lumiInfos, prefix_sum_lumi_size)
make_lumi_summary = lumi_summary_maker(lumiInfos, prefix_sum_lumi_size,
key, lumi_sum_length,
schema_for_algorithms)
return {
"algorithms":
......
###############################################################################
# (c) Copyright 2018-2022 CERN for the benefit of the LHCb Collaboration #
###############################################################################
from sys import argv, exit
from random import Random
import math
import copy
import json
"""Script to generate a LumiSummary bank layout for a set of counters specified in
the input files. The script produces a JSON representation of the layout, which is used
to encode and decode the luminosity summary counters. If the PyConf module is available
the encoding key for the schema will also be calculated.
See the --help option for a full list of options.
This helper script produces an optimised JSON representation for a given collection of luminosity counters.
Input files should contain one counter per line formatted as "CounterName MaxEntries",
where the name may not contain white space and MaxEntries is the maximum value that the counter should be expected to contain.
As multiple counters may be packed into a single 32-bit integer and counters may not overrun between two integers,
the packing order of the counters must be optimised.
The script initially implements a "first-fit decreasing" algorithm where counters are considered in descending size order
and placed into the first available slot (adding an extra int if necessary).
The script attempts to further optimise the schema by randomly removing a subset of the smaller counters (<16bits)
and reinserting then in a random order. Such "mutations" are retained only if they reduce the number of ints required to store the counters.
A schema where all of the words contain a 16-bit or larger counter cannot be further optimised so the second stage is skipped in such cases.
The mutation is repeated a configurable number of times until a configurable packing efficiency is achieved.
An example input file follows:
EXAMPLE
T0Low 0xffffffff
T0High 0xffffffff
BCIDLow 0xffffffff
BCIDHigh 0x3fff
BXType 3
GEC 1
VeloTracks 1913
VeloVertices 33
SciFiClustersS1M45 765
SciFiClustersS2M45 805
SciFiClustersS3M45 1405
SciFiClusters 7650
SciFiClustersS2M123 7590
SciFiClustersS3M123 7890
ECalET 1072742
ECalEInnerTop 3797317
ECalEMiddleTop 1478032
ECalEOuterTop 1192952
ECalEInnerBottom 4026243
ECalEMiddleBottom 1492195
ECalEOuterBottom 1384124
MuonHitsM2R1 696
MuonHitsM2R2 593
MuonHitsM2R3 263
MuonHitsM2R4 200
MuonHitsM3R1 478
MuonHitsM3R2 212
MuonHitsM3R3 161
MuonHitsM3R4 102
MuonHitsM4R1 134
MuonHitsM4R2 108
MuonHitsM4R3 409
MuonHitsM4R4 227
EOF
"""
lumi_rand = Random("lumi_schema_generator")
class Counter:
"""A single lumi counter"""
def __init__(self, name, maxVal):
self.name = name
self.size = math.floor(math.log2(maxVal)) + 1
self.offset = 0
def __lt__(self, other):
return self.size < other.size
class Bucket:
"""A single 32-bit 'bucket' for storing counters"""
def __init__(self, pos):
self.bitsRemaining = 32
self.pos = pos
self.counters = []
def addCounter(self, counter):
"""Put a counter in the bucket"""
if counter.size > self.bitsRemaining:
return False
counter.offset = 32 - self.bitsRemaining
self.bitsRemaining -= counter.size
self.counters.append(counter)
return True
class LumiSchemaGenerator:
"""Class to produce a JSON representation of a LumiSummary bank layout.
Counter names are associated with a size and an offset within the bank.
Multiple counters may be packed into a single 32-bit integer, however,
a single counter may not span more than one integer.
"""
def __init__(self, inputs=[]):
"""Optionally, provide counters as the argument "input" in the form [(counterName1, MAXENTRIES1), (counterName2, MAXENTRIES2) ... ]
"""
self.inputs = [Counter("encodingKey", 0xffffffff)]
for name, maxEntry in inputs:
self.inputs.append(Counter(name, maxEntry))
self.buckets = []
self.size = 0
self.sumSizes = 0
def readInput(self, inputFileName):
"""Append the contents of the input file to the list of requested counters.
Input files should be formatted as follows:
counterName1 MAXENTRIES1
counterName2 MAXENTRIES2
...
where MAXENTRIES is the maximum value that a given counter may be required to store.
Counter names may not contain whitespace. An "encodingKey" counter will be automatically
added to the start of the schema so must not be specified in the input.
If no input file is given then the example from the module docstring is used.
"""
if inputFileName == None:
#Extract the example
lines = open(argv[0]).readlines()
lines = lines[lines.index("EXAMPLE\n") + 1:]
lines = lines[:lines.index("EOF\n")]
else:
lines = open(inputFileName).readlines()
for line in lines:
line = line.split()
if len(line) != 2:
print("Input file, %s, is not in the correct format" %
(inputFileName))
print(
"Expect lines of counter name and maximum value separated by white space"
)
print("Offending line: %s" % " ".join(line))
exit()
(name, maxEntry) = line
self.inputs.append(Counter(name, int(maxEntry, 0)))
print("Found %d counters in %s" % (len(lines), inputFileName))
def processWithoutOptimisation(self):
"""Pack counters into 32-bit bins sequentially without running any optimisation.
If a counter does not fit into the current last bin a new bin is appended.
To use this functionality set the --no-opt command line option.
"""
self.buckets = []
self.size = 0
self.sumSizes = 0
self.nInputs = len(self.inputs)
self.pack(False)
def process(self, mutationAttempts=10, stopThreshold=100.):
"""Pack requested counters into 32-bit bins according to 'first-fit decreasing' procedure.
Counters are sorted in order of descending size and packed into the first available 32-bit 'bucket'.
The schema is further optimised by randomly removing a random fraction of those counters that occupy
less than half a bucket and attempting to re-insert them in a random order. If the number of buckets
is reduced then the 'mutated' schema is retained. This procedure is repeated up to the specified
number of attempts or until the efficiency achieved exceeds the specified stopping threshold.
"""
self.buckets = []
self.size = 0
self.sumSizes = 0
self.nInputs = len(self.inputs)
self.inputs.sort(reverse=True)
self.pack()
#if all buckets contain a counter of 16 bits or larger then no optimisation will reduce the number of buckets
runMutationStep = False
for bucket in self.buckets:
if bucket.counters[0].size < 16:
runMutationStep = True
break
if runMutationStep:
for i in range(mutationAttempts):
if 100. * self.sumSizes / self.size >= stopThreshold:
print(
"Packing efficiency of %.1f%% has reached or exceeded %.1f%%"
% (100. * self.sumSizes / self.size, stopThreshold))
print("Stopping mutation")
break
self.mutate(lumi_rand.random())
def pack(self, optimise=True):
for counter in self.inputs:
bucketFound = False
if optimise:
for bucket in self.buckets:
if bucket.addCounter(counter):
bucketFound = True
break
else:
#if packing is not being optimised then only check the last bucket
if len(self.buckets) > 0:
if self.buckets[-1].addCounter(counter):
bucketFound = True
if not bucketFound:
bucket = Bucket(len(self.buckets))
if bucket.addCounter(counter):
self.buckets.append(bucket)
else:
print(
"Counter %s is too large, ensure it fits within 32 bits"
% counter.name)
exit()
self.sumSizes += counter.size
self.inputs = []
bucket = self.buckets[-1]
counter = bucket.counters[-1]
self.sumSizes = math.ceil(self.sumSizes / 32) * 32
self.size = math.ceil(
(32 * bucket.pos + counter.offset + counter.size) / 32) * 32
print("Packed %d counters into %d bytes" % (self.nInputs,
self.size / 8.))
print("Counter packing is %.1f%% efficient" %
(100 * self.sumSizes / self.size))
def mutate(self, prob):
print("Mutating with a %.1f%% removal rate" % (prob * 100.))
originalBuckets = copy.deepcopy(self.buckets)
originalSize = self.size
originalSumSizes = self.sumSizes
#randomly select counters to re-insert - counters of 16 bits or larger are left intact as removing them is analogous to simply reordering buckets
for bucket in self.buckets:
if bucket.bitsRemaining > 0:
for counter in list(bucket.counters):
if counter.size < 16 and lumi_rand.random() < prob:
bucket.counters.remove(counter)
bucket.bitsRemaining += counter.size
self.sumSizes -= counter.size
self.inputs.append(counter)
#remove any buckets that are now empty
runMutation = False
for bucket in list(self.buckets):
if len(bucket.counters) == 0:
self.buckets.remove(bucket)
runMutation = True
#size can only be reduced if a bucket is removed so skip if no empty buckets
if runMutation:
if len(self.inputs) != 0:
#rerun the packing with a random ordering
print("repacking %d counters" % len(self.inputs))
lumi_rand.shuffle(self.inputs)
self.pack()
if originalSize <= self.size:
print("No improvement, reverting mutation")
self.buckets = originalBuckets
self.size = originalSize
self.sumSizes = originalSumSizes
else:
print("Improvement found, retaining mutation")
else:
self.inputs = []
print("No improvement possible, skipping mutation")
self.buckets = originalBuckets
self.size = originalSize
self.sumSizes = originalSumSizes
def getJSON(self):
"""Return a JSON representation of the lumi counter scheme."""
mk_counter = lambda b,c : { 'name':c.name, 'offset' : 32 * b.pos + c.offset, 'size': c.size }
counters = [
mk_counter(bucket, counter) for bucket in self.buckets
for counter in bucket.counters
]
return {'version': 0, 'size': int(self.size / 8), 'counters': counters}
def printJSON(self):
"""Print a JSON representation of the lumi counter scheme."""
print("JSON representation:")
schema = json.dumps(self.getJSON())
print(schema)
outputName = "output.json"
try:
from PyConf.filecontent_metadata import _get_hash_for_text
schemaKey = _get_hash_for_text(schema)[:8]
print("Encoding key: %s" % (schemaKey))
outputName = schemaKey + ".json"
except ModuleNotFoundError:
print(
"PyConf module unavailable: schema encoding key cannot be calculated"
)
print("Writing %s" % (outputName))
f = open(outputName, "w")
f.write(schema)
f.close()
def printHeaderFile(self):
"""Print enum values for the legacy header interface."""
for b in self.buckets:
for c in b.counters:
print("%sOffset = %d," % (c.name, 32 * b.pos + c.offset))
print("%sSize = %d," % (c.name, c.size))
print("TotalSize = %d" % (self.size))
def printPacking(self):
"""Print a schematic layout of the counter packing within bins."""
print("Counter packing:")
for bucket in self.buckets:
c = '0'
for counter in bucket.counters:
print(c * counter.size, end="")
c = chr(ord(c) + 1)
print("X" * bucket.bitsRemaining)
if __name__ == "__main__":
import argparse
parser = argparse.ArgumentParser()
parser.add_argument(
"--input",
default=None,
help="Input file giving counters and maximum values")
parser.add_argument(
"--no-opt",
metavar="0/1",
nargs='?',
const=True,
default=False,
help="Don't perform any optimisation")
parser.add_argument(
"--mutations",
metavar="N",
type=int,
default=20,
help=
"Number of mutations to perform in the second phase of optimisation")
parser.add_argument(
"--stop-threshold",
metavar="PC",
type=float,
default=100.,
help=
"Packing efficiency at which further mutation attempts should be skipped"
)
parser.add_argument(
"--write-header",
metavar="0/1",
nargs='?',
const=True,
default=False,
help="Generate a header file for the legacy enum interface")
args = parser.parse_args()
l = LumiSchemaGenerator()
l.readInput(args.input)
if args.no_opt:
l.processWithoutOptimisation()
else:
l.process(args.mutations, args.stop_threshold)
l.printPacking()
l.printJSON()
if args.write_header:
l.printHeaderFile()
......@@ -10,7 +10,6 @@
\*****************************************************************************/
#pragma once
#include "LumiSummaryOffsets.h"
#include "MuonDefinitions.cuh"
namespace Lumi {
......@@ -55,9 +54,10 @@ namespace Lumi {
static constexpr unsigned n_plume_channels = 32u;
static constexpr unsigned n_plume_lumi_channels = 22u;
static constexpr unsigned n_basic_counters = 6u;
static constexpr unsigned n_velo_counters = 1u;
static constexpr unsigned n_pv_counters = 1u;
static constexpr unsigned n_SciFi_counters = 6u;
static constexpr unsigned n_scifi_counters = 6u;
static constexpr unsigned n_calo_counters = 7u;
static constexpr unsigned n_muon_counters = 12u;
static constexpr unsigned n_plume_counters = 3u;
......@@ -65,13 +65,43 @@ namespace Lumi {
// number of sub info, used for info aggregating in make_lumi_summary
static constexpr unsigned n_sub_infos = 6u;
// give the length of a lumi summary in unsigned
static constexpr unsigned lumi_length = LHCb::LumiSummaryOffsets::V2::TotalSize / 8u / sizeof(unsigned);
const std::array<std::string, n_basic_counters> basic_counter_names =
{"T0Low", "T0High", "BCIDLow", "BCIDHigh", "BXType", "GEC"};
const std::array<std::string, n_velo_counters> velo_counter_names = {"VeloTracks"};
const std::array<std::string, n_pv_counters> pv_counter_names = {"VeloVertices"};
const std::array<std::string, n_scifi_counters> scifi_counter_names = {"SciFiClusters",
"SciFiClustersS2M123",
"SciFiClustersS3M123",
"SciFiClustersS1M45",
"SciFiClustersS2M45",
"SciFiClustersS3M45"};
const std::array<std::string, n_calo_counters> calo_counter_names = {"ECalET",
"ECalEOuterTop",
"ECalEMiddleTop",
"ECalEInnerTop",
"ECalEOuterBottom",
"ECalEMiddleBottom",
"ECalEInnerBottom"};
const std::array<std::string, n_muon_counters> muon_counter_names = {"MuonHitsM2R1",
"MuonHitsM2R2",
"MuonHitsM2R3",
"MuonHitsM2R4",
"MuonHitsM3R1",
"MuonHitsM3R2",
"MuonHitsM3R3",
"MuonHitsM3R4",
"MuonHitsM4R1",
"MuonHitsM4R2",
"MuonHitsM4R3",
"MuonHitsM4R4"};
const std::array<std::string, n_plume_counters> plume_counter_names = {"PlumeAvgLumiADC",
"PlumeLumiOverthrLow",
"PlumeLumiOverthrHigh"};
} // namespace Constants
struct LumiInfo {
LHCb::LumiSummaryOffsets::V2::counterOffsets size;
LHCb::LumiSummaryOffsets::V2::counterOffsets offset;
unsigned size;
unsigned offset;
unsigned value;
};
} // namespace Lumi
/*****************************************************************************\
* (c) Copyright 2022 CERN for the benefit of the LHCb Collaboration *
* *
* This software is distributed under the terms of the Apache License *
* version 2 (Apache-2.0), copied verbatim in the file "COPYING". *
* *
* In applying this licence, CERN does not waive the privileges and immunities *
* granted to it by virtue of its status as an Intergovernmental Organization *
* or submit itself to any jurisdiction. *
\*****************************************************************************/
#pragma once
#include <iostream>
namespace LHCb {
namespace LumiSummaryOffsets {
namespace V2 {
enum counterOffsets : unsigned {
// All values are in bits - the required size of the array may be determined
// by dividing the largest offset by 8*sizeof(unsigned), i.e. 32, and rounding up.
// Fields must be contained within a single element of the array, e.g. an
// offset of 24 would allow for a maximum size of 8.
// The encoding key stores a hash of the JSON representation of all counter offsets and sizes.
// If a non-default encoding key is present, then all other counters should be decoded using
// the scheme stored in the ANN for that key rather than the default values in this header.
encodingKeyOffset = 0,
encodingKeySize = 32,
T0LowOffset = 32,
T0LowSize = 32,
T0HighOffset = 64,
T0HighSize = 32,
BCIDLowOffset = 96,
BCIDLowSize = 32,
PlumeLumiOverthrLowOffset = 128,
PlumeLumiOverthrLowSize = 22,
MuonHitsM3R2Offset = 150,
MuonHitsM3R2Size = 10,
PlumeLumiOverthrHighOffset = 160,
PlumeLumiOverthrHighSize = 22,
MuonHitsM4R1Offset = 182,
MuonHitsM4R1Size = 10,
SciFiClustersS3M45Offset = 192,
SciFiClustersS3M45Size = 16,
SciFiClustersOffset = 208,
SciFiClustersSize = 16,
SciFiClustersS2M123Offset = 224,
SciFiClustersS2M123Size = 16,
SciFiClustersS3M123Offset = 240,
SciFiClustersS3M123Size = 16,
ECalETOffset = 256,
ECalETSize = 16,
ECalEInnerTopOffset = 272,
ECalEInnerTopSize = 16,
ECalEMiddleTopOffset = 288,
ECalEMiddleTopSize = 16,
ECalEOuterTopOffset = 304,
ECalEOuterTopSize = 16,
ECalEInnerBottomOffset = 320,
ECalEInnerBottomSize = 16,
ECalEMiddleBottomOffset = 336,
ECalEMiddleBottomSize = 16,
ECalEOuterBottomOffset = 352,
ECalEOuterBottomSize = 16,
MuonHitsM2R1Offset = 368,
MuonHitsM2R1Size = 16,
MuonHitsM2R2Offset = 384,
MuonHitsM2R2Size = 16,
MuonHitsM2R3Offset = 400,
MuonHitsM2R3Size = 16,
VeloTracksOffset = 416,
VeloTracksSize = 15,
BCIDHighOffset = 431,
BCIDHighSize = 14,
BXTypeOffset = 445,
BXTypeSize = 2,
GECOffset = 447,
GECSize = 1,
SciFiClustersS1M45Offset = 448,
SciFiClustersS1M45Size = 13,
SciFiClustersS2M45Offset = 461,
SciFiClustersS2M45Size = 13,
VeloVerticesOffset = 474,
VeloVerticesSize = 6,
PlumeAvgLumiADCOffset = 480,
PlumeAvgLumiADCSize = 12,
MuonHitsM2R4Offset = 492,
MuonHitsM2R4Size = 11,
MuonHitsM3R1Offset = 512,
MuonHitsM3R1Size = 11,
MuonHitsM3R3Offset = 523,
MuonHitsM3R3Size = 11,
MuonHitsM4R4Offset = 534,
MuonHitsM4R4Size = 10,
MuonHitsM3R4Offset = 544,
MuonHitsM3R4Size = 11,
MuonHitsM4R2Offset = 555,
MuonHitsM4R2Size = 11,
MuonHitsM4R3Offset = 576,
MuonHitsM4R3Size = 11,
TotalSize = 608
}; // enum CounterOffsets
} // namespace V2
} // namespace LumiSummaryOffsets
} // namespace LHCb
......@@ -20,6 +20,7 @@ namespace calc_lumi_sum_size {
DEVICE_OUTPUT(dev_lumi_sum_sizes_t, unsigned) dev_lumi_sum_sizes;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(line_index_t, "line_index", "index of lumi line", unsigned) line_index;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
}; // struct Parameters
__global__ void calc_lumi_sum_size(Parameters, const unsigned number_of_events);
......@@ -36,5 +37,6 @@ namespace calc_lumi_sum_size {
private:
Property<block_dim_t> m_block_dim {this, {{128, 1, 1}}};
Property<line_index_t> m_line_index {this, 0};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
}; // struct calc_lumi_sum_size_t
} // namespace calc_lumi_sum_size
......@@ -27,6 +27,18 @@ namespace calo_lumi_counters {
DEVICE_INPUT(dev_ecal_digits_offsets_t, unsigned) dev_ecal_digits_offsets;
DEVICE_OUTPUT(dev_lumi_infos_t, Lumi::LumiInfo) dev_lumi_infos;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
PROPERTY(
lumi_counter_schema_t,
"lumi_counter_schema",
"schema for lumi counters",
std::map<std::string, std::pair<unsigned, unsigned>>);
PROPERTY(
calo_offsets_and_sizes_t,
"calo_offsets_and_sizes",
"offsets and sizes in bits for the calo counters",
std::array<unsigned, 2 * Lumi::Constants::n_calo_counters>)
calo_offsets_and_sizes;
}; // struct Parameters
__global__ void calo_lumi_counters(Parameters, const unsigned number_of_events, const char* raw_ecal_geometry);
......@@ -35,6 +47,8 @@ namespace calo_lumi_counters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
......@@ -43,6 +57,11 @@ namespace calo_lumi_counters {
private:
Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
Property<lumi_counter_schema_t> m_lumi_counter_schema {this, {}};
Property<calo_offsets_and_sizes_t> m_calo_offsets_and_sizes {
this,
{{0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u}}};
}; // struct calo_lumi_counters_t
} // namespace calo_lumi_counters
/*****************************************************************************\
* (c) Copyright 2022 CERN for the benefit of the LHCb Collaboration *
* *
* This software is distributed under the terms of the Apache License *
* version 2 (Apache-2.0), copied verbatim in the file "COPYING". *
* *
* In applying this licence, CERN does not waive the privileges and immunities *
* granted to it by virtue of its status as an Intergovernmental Organization *
* or submit itself to any jurisdiction. *
\*****************************************************************************/
#pragma once
#include <LumiDefinitions.cuh>
inline __device__ void
fillLumiInfo(Lumi::LumiInfo& info, const unsigned offset, const unsigned size, const unsigned value)
{
info.size = size;
info.offset = offset;
info.value = value;
}
......@@ -14,7 +14,6 @@
#include "AlgorithmTypes.cuh"
#include "GenericContainerContracts.h"
#include "LumiSummaryOffsets.h"
#include <LumiDefinitions.cuh>
#include "ODINBank.cuh"
......@@ -36,6 +35,18 @@ namespace make_lumi_summary {
HOST_OUTPUT(host_lumi_summary_offsets_t, unsigned) host_lumi_summary_offsets;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(encoding_key_t, "encoding_key", "encoding key", unsigned) key;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
PROPERTY(
lumi_counter_schema_t,
"lumi_counter_schema",
"schema for lumi counters",
std::map<std::string, std::pair<unsigned, unsigned>>);
PROPERTY(
basic_offsets_and_sizes_t,
"basic_offsets_and_sizes",
"offsets and sizes in bits for the ODIN and GEC counters",
std::array<unsigned, 2 * Lumi::Constants::n_basic_counters>)
basic_offsets_and_sizes;
}; // struct Parameters
__global__ void make_lumi_summary(
......@@ -46,15 +57,13 @@ namespace make_lumi_summary {
std::array<unsigned, Lumi::Constants::n_sub_infos> spanSize,
const unsigned size_of_aggregate);
__device__ void setField(
LHCb::LumiSummaryOffsets::V2::counterOffsets offset,
LHCb::LumiSummaryOffsets::V2::counterOffsets size,
unsigned* target,
unsigned value);
__device__ void setField(unsigned offset, unsigned size, unsigned* target, unsigned value);
struct make_lumi_summary_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
......@@ -64,5 +73,9 @@ namespace make_lumi_summary {
private:
Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
Property<encoding_key_t> m_key {this, 0};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
Property<lumi_counter_schema_t> m_lumi_counter_schema {this, {}};
Property<basic_offsets_and_sizes_t> m_basic_offsets_and_sizes {this,
{{0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u}}};
}; // struct make_lumi_summary_t
} // namespace make_lumi_summary
......@@ -26,6 +26,18 @@ namespace muon_lumi_counters {
DEVICE_INPUT(dev_storage_station_region_quarter_offsets_t, unsigned) dev_storage_station_region_quarter_offsets;
DEVICE_OUTPUT(dev_lumi_infos_t, Lumi::LumiInfo) dev_lumi_infos;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
PROPERTY(
lumi_counter_schema_t,
"lumi_counter_schema",
"schema for lumi counters",
std::map<std::string, std::pair<unsigned, unsigned>>);
PROPERTY(
muon_offsets_and_sizes_t,
"muon_offsets_and_sizes",
"offsets and sizes in bits for the muon counters",
std::array<unsigned, 2 * Lumi::Constants::n_muon_counters>)
muon_offsets_and_sizes;
}; // struct Parameters
__global__ void muon_lumi_counters(Parameters, const unsigned number_of_events);
......@@ -33,6 +45,8 @@ namespace muon_lumi_counters {
struct muon_lumi_counters_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
......@@ -41,5 +55,10 @@ namespace muon_lumi_counters {
private:
Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
Property<lumi_counter_schema_t> m_lumi_counter_schema {this, {}};
Property<muon_offsets_and_sizes_t> m_muon_offsets_and_sizes {
this,
{{0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u}}};
}; // struct muon_lumi_counters_t
} // namespace muon_lumi_counters
......@@ -27,6 +27,18 @@ namespace pv_lumi_counters {
DEVICE_INPUT(dev_number_of_pvs_t, unsigned) dev_number_of_pvs;
DEVICE_OUTPUT(dev_lumi_infos_t, Lumi::LumiInfo) dev_lumi_infos;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
PROPERTY(
lumi_counter_schema_t,
"lumi_counter_schema",
"schema for lumi counters",
std::map<std::string, std::pair<unsigned, unsigned>>);
PROPERTY(
pv_offsets_and_sizes_t,
"pv_offsets_and_sizes",
"offsets and sizes in bits for the PV counters",
std::array<unsigned, 2 * Lumi::Constants::n_pv_counters>)
pv_offsets_and_sizes;
}; // struct Parameters
__global__ void pv_lumi_counters(Parameters, const unsigned number_of_events);
......@@ -34,6 +46,8 @@ namespace pv_lumi_counters {
struct pv_lumi_counters_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
......@@ -42,5 +56,8 @@ namespace pv_lumi_counters {
private:
Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
Property<lumi_counter_schema_t> m_lumi_counter_schema {this, {}};
Property<pv_offsets_and_sizes_t> m_pv_offsets_and_sizes {this, {{0u, 0u}}};
}; // struct pv_lumi_counters_t
} // namespace pv_lumi_counters
......@@ -26,6 +26,18 @@ namespace plume_lumi_counters {
DEVICE_INPUT(dev_plume_t, Plume_) dev_plume;
DEVICE_OUTPUT(dev_lumi_infos_t, Lumi::LumiInfo) dev_lumi_infos;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
PROPERTY(
lumi_counter_schema_t,
"lumi_counter_schema",
"schema for lumi counters",
std::map<std::string, std::pair<unsigned, unsigned>>);
PROPERTY(
plume_offsets_and_sizes_t,
"plume_offsets_and_sizes",
"offsets and sizes in bits for the PV counters",
std::array<unsigned, 2 * Lumi::Constants::n_plume_counters>)
plume_offsets_and_sizes;
}; // struct Parameters
__global__ void plume_lumi_counters(Parameters, const unsigned number_of_events);
......@@ -33,6 +45,8 @@ namespace plume_lumi_counters {
struct plume_lumi_counters_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
......@@ -41,5 +55,8 @@ namespace plume_lumi_counters {
private:
Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
Property<lumi_counter_schema_t> m_lumi_counter_schema {this, {}};
Property<plume_offsets_and_sizes_t> m_plume_offsets_and_sizes {this, {{0u, 0u, 0u, 0u, 0u, 0u}}};
}; // struct plume_lumi_counters_t
} // namespace plume_lumi_counters
......@@ -25,6 +25,18 @@ namespace scifi_lumi_counters {
DEVICE_INPUT(dev_scifi_hits_t, char) dev_scifi_hits;
DEVICE_OUTPUT(dev_lumi_infos_t, Lumi::LumiInfo) dev_lumi_infos;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
PROPERTY(
lumi_counter_schema_t,
"lumi_counter_schema",
"schema for lumi counters",
std::map<std::string, std::pair<unsigned, unsigned>>);
PROPERTY(
scifi_offsets_and_sizes_t,
"scifi_offsets_and_sizes",
"offsets and sizes in bits for the SciFi counters",
std::array<unsigned, 2 * Lumi::Constants::n_scifi_counters>)
scifi_offsets_and_sizes;
}; // struct Parameters
__global__ void scifi_lumi_counters(Parameters, const unsigned number_of_events, const char* scifi_geometry);
......@@ -32,6 +44,8 @@ namespace scifi_lumi_counters {
struct scifi_lumi_counters_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
......@@ -40,5 +54,9 @@ namespace scifi_lumi_counters {
private:
Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
Property<lumi_counter_schema_t> m_lumi_counter_schema {this, {}};
Property<scifi_offsets_and_sizes_t> m_scifi_offsets_and_sizes {this,
{{0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u}}};
}; // struct scifi_lumi_counters_t
} // namespace scifi_lumi_counters
......@@ -27,6 +27,18 @@ namespace velo_lumi_counters {
DEVICE_INPUT(dev_offsets_all_velo_tracks_t, unsigned) dev_offsets_all_velo_tracks;
DEVICE_OUTPUT(dev_lumi_infos_t, Lumi::LumiInfo) dev_lumi_infos;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(lumi_sum_length_t, "lumi_sum_length", "LumiSummary length", unsigned) lumi_sum_length;
PROPERTY(
lumi_counter_schema_t,
"lumi_counter_schema",
"schema for lumi counters",
std::map<std::string, std::pair<unsigned, unsigned>>);
PROPERTY(
velo_offsets_and_sizes_t,
"velo_offsets_and_sizes",
"offsets and sizes in bits for the VELO counters",
std::array<unsigned, 2 * Lumi::Constants::n_velo_counters>)
velo_offsets_and_sizes;
}; // struct Parameters
__global__ void velo_lumi_counters(Parameters, const unsigned number_of_events);
......@@ -34,6 +46,8 @@ namespace velo_lumi_counters {
struct velo_lumi_counters_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
......@@ -42,5 +56,8 @@ namespace velo_lumi_counters {
private:
Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
Property<lumi_sum_length_t> m_lumi_sum_length {this, 0u};
Property<lumi_counter_schema_t> m_lumi_counter_schema {this, {}};
Property<velo_offsets_and_sizes_t> m_velo_offsets_and_sizes {this, {{0u, 0u}}};
}; // struct velo_lumi_counters_t
} // namespace velo_lumi_counters
......@@ -12,7 +12,6 @@
#include "SelectionsEventModel.cuh"
#include "LumiDefinitions.cuh"
#include "LumiSummaryOffsets.h"
INSTANTIATE_ALGORITHM(calc_lumi_sum_size::calc_lumi_sum_size_t)
......@@ -49,8 +48,6 @@ __global__ void calc_lumi_sum_size::calc_lumi_sum_size(
const auto sel_span = selections.get_span(parameters.line_index, event_number);
if (sel_span.empty() || !sel_span[0]) continue;
// the unit of LHCb::LumiSummaryOffsets::V2::TotalSize is bit
// convert it to be unsigned int
parameters.dev_lumi_sum_sizes[event_number] = Lumi::Constants::lumi_length;
parameters.dev_lumi_sum_sizes[event_number] = parameters.lumi_sum_length;
}
}
......@@ -9,7 +9,7 @@
* or submit itself to any jurisdiction. *
\*****************************************************************************/
#include "CaloLumiCounters.cuh"
#include "LumiSummaryOffsets.h"
#include "LumiCommon.cuh"
#include "CaloGeometry.cuh"
......@@ -23,7 +23,27 @@ void calo_lumi_counters::calo_lumi_counters_t::set_arguments_size(
// the total size of output info is proportional to the lumi summaries
set_size<dev_lumi_infos_t>(
arguments,
Lumi::Constants::n_calo_counters * first<host_lumi_summaries_size_t>(arguments) / Lumi::Constants::lumi_length);
Lumi::Constants::n_calo_counters * first<host_lumi_summaries_size_t>(arguments) / property<lumi_sum_length_t>());
}
void calo_lumi_counters::calo_lumi_counters_t::init()
{
std::map<std::string, std::pair<unsigned, unsigned>> schema = property<lumi_counter_schema_t>();
std::array<unsigned, 2 * Lumi::Constants::n_calo_counters> calo_offsets_and_sizes =
property<calo_offsets_and_sizes_t>();
unsigned c_idx(0u);
for (auto counter_name : Lumi::Constants::calo_counter_names) {
if (schema.find(counter_name) == schema.end()) {
std::cout << "LumiSummary schema does not use " << counter_name << std::endl;
}
else {
calo_offsets_and_sizes[2 * c_idx] = schema[counter_name].first;
calo_offsets_and_sizes[2 * c_idx + 1] = schema[counter_name].second;
}
++c_idx;
}
set_property_value<calo_offsets_and_sizes_t>(calo_offsets_and_sizes);
}
void calo_lumi_counters::calo_lumi_counters_t::operator()(
......@@ -55,8 +75,8 @@ __global__ void calo_lumi_counters::calo_lumi_counters(
const unsigned digits_offset = parameters.dev_ecal_digits_offsets[event_number];
const unsigned n_digits = parameters.dev_ecal_digits_offsets[event_number + 1] - digits_offset;
auto const* digits = parameters.dev_ecal_digits + digits_offset;
float sum_et = 0.f;
std::array<float, 6> sum_et_area = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
// sumET followed by ET for each region
std::array<float, Lumi::Constants::n_calo_counters> E_vals = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
for (unsigned digit_index = 0u; digit_index < n_digits; ++digit_index) {
if (!digits[digit_index].is_valid()) continue;
......@@ -68,57 +88,25 @@ __global__ void calo_lumi_counters::calo_lumi_counters(
auto e = ecal_geometry.getE(digit_index, digits[digit_index].adc);
auto sin_theta = sqrtf((x * x + y * y) / (x * x + y * y + z * z));
sum_et += e * sin_theta;
E_vals[0] += e * sin_theta;
auto const area = ecal_geometry.getECALArea(digit_index);
if (y > 0.f) {
sum_et_area[area] += e * sin_theta;
E_vals[1 + area] += e * sin_theta;
}
else {
sum_et_area[3 + area] += e * sin_theta;
E_vals[4 + area] += e * sin_theta;
}
}
unsigned info_offset = Lumi::Constants::n_calo_counters * lumi_sum_offset / Lumi::Constants::lumi_length;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::ECalETOffset;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::ECalETSize;
parameters.dev_lumi_infos[info_offset].value = sum_et;
// Outer Top
++info_offset;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::ECalEOuterTopOffset;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::ECalEOuterTopSize;
parameters.dev_lumi_infos[info_offset].value = sum_et_area[0];
// Middle Top
++info_offset;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::ECalEMiddleTopOffset;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::ECalEMiddleTopSize;
parameters.dev_lumi_infos[info_offset].value = sum_et_area[1];
// Inner Top
++info_offset;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::ECalEInnerTopOffset;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::ECalEInnerTopSize;
parameters.dev_lumi_infos[info_offset].value = sum_et_area[2];
// Outer Bottom
++info_offset;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::ECalEOuterBottomOffset;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::ECalEOuterBottomSize;
parameters.dev_lumi_infos[info_offset].value = sum_et_area[3];
// Middle Bottom
++info_offset;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::ECalEMiddleBottomOffset;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::ECalEMiddleBottomSize;
parameters.dev_lumi_infos[info_offset].value = sum_et_area[4];
// Inner Bottom
++info_offset;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::ECalEInnerBottomOffset;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::ECalEInnerBottomSize;
parameters.dev_lumi_infos[info_offset].value = sum_et_area[5];
unsigned info_offset = Lumi::Constants::n_calo_counters * lumi_sum_offset / parameters.lumi_sum_length;
for (unsigned i = 0; i < Lumi::Constants::n_calo_counters; ++i) {
fillLumiInfo(
parameters.dev_lumi_infos[info_offset + i],
parameters.calo_offsets_and_sizes.get()[2 * i],
parameters.calo_offsets_and_sizes.get()[2 * i + 1],
E_vals[i]);
}
}
}
......@@ -9,7 +9,6 @@
* or submit itself to any jurisdiction. *
\*****************************************************************************/
#include "MakeLumiSummary.cuh"
#include "LumiSummaryOffsets.h"
#include "SelectionsEventModel.cuh"
#include "Event/ODIN.h"
......@@ -26,6 +25,26 @@ void make_lumi_summary::make_lumi_summary_t::set_arguments_size(
set_size<dev_lumi_summaries_t>(arguments, first<host_lumi_summaries_size_t>(arguments));
}
void make_lumi_summary::make_lumi_summary_t::init()
{
std::map<std::string, std::pair<unsigned, unsigned>> schema = property<lumi_counter_schema_t>();
std::array<unsigned, 2 * Lumi::Constants::n_basic_counters> basic_offsets_and_sizes =
property<basic_offsets_and_sizes_t>();
unsigned c_idx(0u);
for (auto counter_name : Lumi::Constants::basic_counter_names) {
if (schema.find(counter_name) == schema.end()) {
std::cout << "LumiSummary schema does not use " << counter_name << std::endl;
}
else {
basic_offsets_and_sizes[2 * c_idx] = schema[counter_name].first;
basic_offsets_and_sizes[2 * c_idx + 1] = schema[counter_name].second;
}
++c_idx;
}
set_property_value<basic_offsets_and_sizes_t>(basic_offsets_and_sizes);
}
void make_lumi_summary::make_lumi_summary_t::operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions&,
......@@ -45,7 +64,7 @@ void make_lumi_summary::make_lumi_summary_t::operator()(
std::array<unsigned, Lumi::Constants::n_sub_infos> infoSize = {
std::min(Lumi::Constants::n_velo_counters, static_cast<unsigned>(size<dev_velo_info_t>(arguments))),
std::min(Lumi::Constants::n_pv_counters, static_cast<unsigned>(size<dev_pv_info_t>(arguments))),
std::min(Lumi::Constants::n_SciFi_counters, static_cast<unsigned>(size<dev_scifi_info_t>(arguments))),
std::min(Lumi::Constants::n_scifi_counters, static_cast<unsigned>(size<dev_scifi_info_t>(arguments))),
std::min(Lumi::Constants::n_muon_counters, static_cast<unsigned>(size<dev_muon_info_t>(arguments))),
std::min(Lumi::Constants::n_calo_counters, static_cast<unsigned>(size<dev_calo_info_t>(arguments))),
std::min(Lumi::Constants::n_plume_counters, static_cast<unsigned>(size<dev_plume_info_t>(arguments)))};
......@@ -75,11 +94,7 @@ void make_lumi_summary::make_lumi_summary_t::operator()(
Allen::copy_async<host_lumi_summary_offsets_t, dev_lumi_summary_offsets_t>(arguments, context);
}
__device__ void make_lumi_summary::setField(
LHCb::LumiSummaryOffsets::V2::counterOffsets offset,
LHCb::LumiSummaryOffsets::V2::counterOffsets size,
unsigned* target,
unsigned value)
__device__ void make_lumi_summary::setField(unsigned offset, unsigned size, unsigned* target, unsigned value)
{
// Check value fits within size bits
if (size < (8 * sizeof(unsigned)) && value >= (1u << size)) {
......@@ -124,49 +139,55 @@ __global__ void make_lumi_summary::make_lumi_summary(
uint64_t t0 = static_cast<uint64_t>(odin.gpsTime()) - new_bcid * 1000 / 40078;
// event time
setField(
LHCb::LumiSummaryOffsets::V2::T0LowOffset,
LHCb::LumiSummaryOffsets::V2::T0LowSize,
parameters.basic_offsets_and_sizes.get()[0],
parameters.basic_offsets_and_sizes.get()[1],
lumi_summary,
static_cast<unsigned>(t0 & 0xffffffff));
setField(
LHCb::LumiSummaryOffsets::V2::T0HighOffset,
LHCb::LumiSummaryOffsets::V2::T0HighSize,
parameters.basic_offsets_and_sizes.get()[2],
parameters.basic_offsets_and_sizes.get()[3],
lumi_summary,
static_cast<unsigned>(t0 >> 32));
// gps time offset
setField(
LHCb::LumiSummaryOffsets::V2::BCIDLowOffset,
LHCb::LumiSummaryOffsets::V2::BCIDLowSize,
parameters.basic_offsets_and_sizes.get()[4],
parameters.basic_offsets_and_sizes.get()[5],
lumi_summary,
static_cast<unsigned>(new_bcid & 0xffffffff));
setField(
LHCb::LumiSummaryOffsets::V2::BCIDHighOffset,
LHCb::LumiSummaryOffsets::V2::BCIDHighSize,
parameters.basic_offsets_and_sizes.get()[6],
parameters.basic_offsets_and_sizes.get()[7],
lumi_summary,
static_cast<unsigned>(new_bcid >> 32));
// bunch crossing type
setField(
LHCb::LumiSummaryOffsets::V2::BXTypeOffset,
LHCb::LumiSummaryOffsets::V2::BXTypeSize,
parameters.basic_offsets_and_sizes.get()[8],
parameters.basic_offsets_and_sizes.get()[9],
lumi_summary,
static_cast<unsigned>(odin.bunchCrossingType()));
/// gec counter
bool passedGEC = false;
for (unsigned i = 0; i < number_of_events_passed_gec; ++i) {
if (parameters.dev_event_list[i] == event_number) {
setField(LHCb::LumiSummaryOffsets::V2::GECOffset, LHCb::LumiSummaryOffsets::V2::GECSize, lumi_summary, true);
passedGEC = true;
break;
}
}
setField(
parameters.basic_offsets_and_sizes.get()[10],
parameters.basic_offsets_and_sizes.get()[11],
lumi_summary,
passedGEC);
/// write lumi infos to the summary
for (unsigned i = 0; i < size_of_aggregate; ++i) {
if (infoSize[i] == 0 || lumiInfos[i] == nullptr) continue;
unsigned spanOffset = offset / Lumi::Constants::lumi_length * infoSize[i];
unsigned spanOffset = offset / parameters.lumi_sum_length * infoSize[i];
for (unsigned j = spanOffset;
j < parameters.dev_lumi_summary_offsets[event_number + 1] / Lumi::Constants::lumi_length * infoSize[i];
j < parameters.dev_lumi_summary_offsets[event_number + 1] / parameters.lumi_sum_length * infoSize[i];
++j) {
setField(lumiInfos[i][j].offset, lumiInfos[i][j].size, lumi_summary, lumiInfos[i][j].value);
}
......
......@@ -9,7 +9,7 @@
* or submit itself to any jurisdiction. *
\*****************************************************************************/
#include "MuonLumiCounters.cuh"
#include "LumiSummaryOffsets.h"
#include "LumiCommon.cuh"
INSTANTIATE_ALGORITHM(muon_lumi_counters::muon_lumi_counters_t)
......@@ -21,7 +21,27 @@ void muon_lumi_counters::muon_lumi_counters_t::set_arguments_size(
// the total size of output info is proportional to the lumi summaries
set_size<dev_lumi_infos_t>(
arguments,
Lumi::Constants::n_muon_counters * first<host_lumi_summaries_size_t>(arguments) / Lumi::Constants::lumi_length);
Lumi::Constants::n_muon_counters * first<host_lumi_summaries_size_t>(arguments) / property<lumi_sum_length_t>());
}
void muon_lumi_counters::muon_lumi_counters_t::init()
{
std::map<std::string, std::pair<unsigned, unsigned>> schema = property<lumi_counter_schema_t>();
std::array<unsigned, 2 * Lumi::Constants::n_muon_counters> muon_offsets_and_sizes =
property<muon_offsets_and_sizes_t>();
unsigned c_idx(0u);
for (auto counter_name : Lumi::Constants::muon_counter_names) {
if (schema.find(counter_name) == schema.end()) {
std::cout << "LumiSummary schema does not use " << counter_name << std::endl;
}
else {
muon_offsets_and_sizes[2 * c_idx] = schema[counter_name].first;
muon_offsets_and_sizes[2 * c_idx + 1] = schema[counter_name].second;
}
++c_idx;
}
set_property_value<muon_offsets_and_sizes_t>(muon_offsets_and_sizes);
}
void muon_lumi_counters::muon_lumi_counters_t::operator()(
......@@ -51,88 +71,28 @@ __global__ void muon_lumi_counters::muon_lumi_counters(
const auto muon_hits_offsets =
parameters.dev_storage_station_region_quarter_offsets + event_number * Lumi::Constants::MuonBankSize;
unsigned muon_info_offset = 12u * lumi_sum_offset / Lumi::Constants::lumi_length;
// M2R1
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R1Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R1Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M2R2] - muon_hits_offsets[Lumi::Constants::M2R1];
// M2R2
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R2Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R2Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M2R3] - muon_hits_offsets[Lumi::Constants::M2R2];
// M2R3
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R3Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R3Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M2R4] - muon_hits_offsets[Lumi::Constants::M2R3];
// M2R4
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R4Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM2R4Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M3R1] - muon_hits_offsets[Lumi::Constants::M2R4];
// M3R1
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R1Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R1Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M3R2] - muon_hits_offsets[Lumi::Constants::M3R1];
// M3R2
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R2Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R2Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M3R3] - muon_hits_offsets[Lumi::Constants::M3R2];
// M3R3
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R3Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R3Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M3R4] - muon_hits_offsets[Lumi::Constants::M3R3];
// M3R4
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R4Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM3R4Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M4R1] - muon_hits_offsets[Lumi::Constants::M3R4];
// M4R1
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R1Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R1Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M4R2] - muon_hits_offsets[Lumi::Constants::M4R1];
// M4R2
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R2Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R2Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M4R3] - muon_hits_offsets[Lumi::Constants::M4R2];
// M4R3
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R3Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R3Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::M4R4] - muon_hits_offsets[Lumi::Constants::M3R3];
// M4R4
++muon_info_offset;
parameters.dev_lumi_infos[muon_info_offset].size = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R4Size;
parameters.dev_lumi_infos[muon_info_offset].offset = LHCb::LumiSummaryOffsets::V2::MuonHitsM4R4Offset;
parameters.dev_lumi_infos[muon_info_offset].value =
muon_hits_offsets[Lumi::Constants::MuonBankSize] - muon_hits_offsets[Lumi::Constants::M4R4];
unsigned info_offset = 12u * lumi_sum_offset / parameters.lumi_sum_length;
std::array<unsigned, Lumi::Constants::n_muon_counters + 1> muon_offsets = {Lumi::Constants::M2R1,
Lumi::Constants::M2R2,
Lumi::Constants::M2R3,
Lumi::Constants::M2R4,
Lumi::Constants::M3R1,
Lumi::Constants::M3R2,
Lumi::Constants::M3R3,
Lumi::Constants::M3R4,
Lumi::Constants::M4R1,
Lumi::Constants::M4R2,
Lumi::Constants::M4R3,
Lumi::Constants::M4R4,
Lumi::Constants::MuonBankSize};
for (unsigned i = 0; i < Lumi::Constants::n_muon_counters; ++i) {
fillLumiInfo(
parameters.dev_lumi_infos[info_offset + i],
parameters.muon_offsets_and_sizes.get()[2 * i],
parameters.muon_offsets_and_sizes.get()[2 * i + 1],
muon_hits_offsets[muon_offsets[i + 1]] - muon_hits_offsets[muon_offsets[i]]);
}
}
}
......@@ -9,7 +9,7 @@
* or submit itself to any jurisdiction. *
\*****************************************************************************/
#include "PVLumiCounters.cuh"
#include "LumiSummaryOffsets.h"
#include "LumiCommon.cuh"
INSTANTIATE_ALGORITHM(pv_lumi_counters::pv_lumi_counters_t)
......@@ -21,7 +21,26 @@ void pv_lumi_counters::pv_lumi_counters_t::set_arguments_size(
// the total size of output info is proportional to the lumi summaries
set_size<dev_lumi_infos_t>(
arguments,
Lumi::Constants::n_pv_counters * first<host_lumi_summaries_size_t>(arguments) / Lumi::Constants::lumi_length);
Lumi::Constants::n_pv_counters * first<host_lumi_summaries_size_t>(arguments) / property<lumi_sum_length_t>());
}
void pv_lumi_counters::pv_lumi_counters_t::init()
{
std::map<std::string, std::pair<unsigned, unsigned>> schema = property<lumi_counter_schema_t>();
std::array<unsigned, 2 * Lumi::Constants::n_pv_counters> pv_offsets_and_sizes = property<pv_offsets_and_sizes_t>();
unsigned c_idx(0u);
for (auto counter_name : Lumi::Constants::pv_counter_names) {
if (schema.find(counter_name) == schema.end()) {
std::cout << "LumiSummary schema does not use " << counter_name << std::endl;
}
else {
pv_offsets_and_sizes[2 * c_idx] = schema[counter_name].first;
pv_offsets_and_sizes[2 * c_idx + 1] = schema[counter_name].second;
}
++c_idx;
}
set_property_value<pv_offsets_and_sizes_t>(pv_offsets_and_sizes);
}
void pv_lumi_counters::pv_lumi_counters_t::operator()(
......@@ -49,9 +68,12 @@ __global__ void pv_lumi_counters::pv_lumi_counters(
if (lumi_sum_offset == parameters.dev_lumi_summary_offsets[event_number + 1]) continue;
// number of PVs
unsigned info_offset = lumi_sum_offset / Lumi::Constants::lumi_length;
parameters.dev_lumi_infos[info_offset].size = LHCb::LumiSummaryOffsets::V2::VeloVerticesSize;
parameters.dev_lumi_infos[info_offset].offset = LHCb::LumiSummaryOffsets::V2::VeloVerticesOffset;
parameters.dev_lumi_infos[info_offset].value = parameters.dev_number_of_pvs[event_number];
unsigned info_offset = lumi_sum_offset / parameters.lumi_sum_length;
fillLumiInfo(
parameters.dev_lumi_infos[info_offset],
parameters.pv_offsets_and_sizes.get()[0],
parameters.pv_offsets_and_sizes.get()[1],
parameters.dev_number_of_pvs[event_number]);
}
}