diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml deleted file mode 100644 index df3668dbc30e1e883f829636a5201df1e73367cc..0000000000000000000000000000000000000000 --- a/.gitlab-ci.yml +++ /dev/null @@ -1,715 +0,0 @@ -stages: - - check-env # Ensures the CI environment is valid - - build # Builds all projects - - run # Runs the tests across architectures - - test # Runs various tests of the software - - publish # Publishes the results of the tests and runs in channels and grafana - -check-env: - stage: check-env - except: - - /.*/@lhcb/Allen - script: - - | - echo "The Allen CI depends on custom GitLab runners and therefore tests" - echo "running on forks will fail. Please create a branch in the main" - echo "repository at https://gitlab.cern.ch/lhcb/Allen/" - - exit 1 - -check-formatting: - stage: build - image: gitlab-registry.cern.ch/lhcb-docker/style-checker - script: - - if [ ! -e .clang-format ] ; then - - curl -o .clang-format "https://gitlab.cern.ch/lhcb-parallelization/Allen/raw/master/.clang-format?inline=false" - - echo '.clang-format' >> .gitignore - - git add .gitignore - - fi - - curl -o lb-format "https://gitlab.cern.ch/lhcb-core/LbDevTools/raw/master/LbDevTools/SourceTools.py?inline=false" - - python lb-format --format-patch apply-formatting.patch origin/master - artifacts: - paths: - - apply-formatting.patch - when: on_failure - expire_in: 1 week - allow_failure: true - -.build_job: &build_gcc_job_def_warnings_as_errors - only: - refs: - - master - - schedules - - web - - merge_requests - stage: build - script: - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=${PREVIOUS_IFS} - - TARGET=${JOB_NAME_SPLIT[0]} - - SEQUENCE=${JOB_NAME_SPLIT[1]} - - BUILD_TYPE=${JOB_NAME_SPLIT[2]} - - ADDITIONAL_OPTIONS="${JOB_NAME_SPLIT[3]}" - - OVERRIDE_CUDA_ARCH_FLAG="-gencode=arch=compute_70,code=sm_70 -gencode=arch=compute_75,code=sm_75" - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - source /cvmfs/sft.cern.ch/lcg/contrib/cuda/11.0RC/x86_64-centos7/setup.sh - - mkdir build_${TARGET} - - cd build_${TARGET} - - cmake -DSTANDALONE=ON -G Ninja -DTREAT_WARNINGS_AS_ERRORS=ON -DTARGET_DEVICE=${TARGET} -DCMAKE_BUILD_TYPE=${BUILD_TYPE} -DSEQUENCE=${SEQUENCE} -DCPU_ARCH=haswell -DOVERRIDE_CUDA_ARCH_FLAG="${OVERRIDE_CUDA_ARCH_FLAG}" "${ADDITIONAL_OPTIONS}" .. - - ninja -j 8 - artifacts: - expire_in: 2 hrs - paths: - - build*/*Allen* - - build*/Sequence.json - - input - retry: 0 - allow_failure: true - -.build_job: &build_clang_job_def_warnings_as_errors - only: - refs: - - master - - schedules - - web - - merge_requests - stage: build - script: - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=${PREVIOUS_IFS} - - TARGET=${JOB_NAME_SPLIT[0]} - - SEQUENCE=${JOB_NAME_SPLIT[1]} - - BUILD_TYPE=${JOB_NAME_SPLIT[2]} - - ADDITIONAL_OPTIONS="${JOB_NAME_SPLIT[3]}" - - OVERRIDE_CUDA_ARCH_FLAG="-gencode=arch=compute_70,code=sm_70 -gencode=arch=compute_75,code=sm_75" - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-clang10-opt - - source /cvmfs/sft.cern.ch/lcg/contrib/cuda/11.0RC/x86_64-centos7/setup.sh - - mkdir build_${TARGET} - - cd build_${TARGET} - - cmake -DSTANDALONE=ON -G Ninja -DTREAT_WARNINGS_AS_ERRORS=ON -DTARGET_DEVICE=${TARGET} -DCMAKE_BUILD_TYPE=${BUILD_TYPE} -DSEQUENCE=${SEQUENCE} -DCPU_ARCH=haswell -DOVERRIDE_CUDA_ARCH_FLAG="${OVERRIDE_CUDA_ARCH_FLAG}" "${ADDITIONAL_OPTIONS}" .. - - ninja -j 8 - artifacts: - expire_in: 2 hrs - paths: - - build*/*Allen* - - build*/Sequence.json - - input - retry: 0 - allow_failure: true - -.build_job: &build_gcc_job_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: build - script: - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=${PREVIOUS_IFS} - - TARGET=${JOB_NAME_SPLIT[0]} - - SEQUENCE=${JOB_NAME_SPLIT[1]} - - BUILD_TYPE=${JOB_NAME_SPLIT[2]} - - ADDITIONAL_OPTIONS="${JOB_NAME_SPLIT[3]}" - - OVERRIDE_CUDA_ARCH_FLAG="-gencode=arch=compute_70,code=sm_70 -gencode=arch=compute_75,code=sm_75" - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - source /cvmfs/sft.cern.ch/lcg/contrib/cuda/11.0RC/x86_64-centos7/setup.sh - - mkdir build_${TARGET} - - cd build_${TARGET} - - cmake -DSTANDALONE=ON -G Ninja -DTARGET_DEVICE=${TARGET} -DCMAKE_BUILD_TYPE=${BUILD_TYPE} -DSEQUENCE=${SEQUENCE} -DCPU_ARCH=haswell -DOVERRIDE_CUDA_ARCH_FLAG="${OVERRIDE_CUDA_ARCH_FLAG}" "${ADDITIONAL_OPTIONS}" .. - - ninja -j 8 - artifacts: - expire_in: 2 hrs - paths: - - build*/*Allen* - - build*/Sequence.json - - input - retry: 1 - -.build_job: &build_clang_job_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: build - script: - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=${PREVIOUS_IFS} - - TARGET=${JOB_NAME_SPLIT[0]} - - SEQUENCE=${JOB_NAME_SPLIT[1]} - - BUILD_TYPE=${JOB_NAME_SPLIT[2]} - - ADDITIONAL_OPTIONS="${JOB_NAME_SPLIT[3]}" - - OVERRIDE_CUDA_ARCH_FLAG="-gencode=arch=compute_70,code=sm_70 -gencode=arch=compute_75,code=sm_75" - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-clang10-opt - - source /cvmfs/sft.cern.ch/lcg/contrib/cuda/11.0RC/x86_64-centos7/setup.sh - - mkdir build_${TARGET} - - cd build_${TARGET} - - cmake -DSTANDALONE=ON -G Ninja -DTARGET_DEVICE=${TARGET} -DCMAKE_BUILD_TYPE=${BUILD_TYPE} -DSEQUENCE=${SEQUENCE} -DCPU_ARCH=haswell -DOVERRIDE_CUDA_ARCH_FLAG="${OVERRIDE_CUDA_ARCH_FLAG}" "${ADDITIONAL_OPTIONS}" .. - - ninja -j 8 - artifacts: - expire_in: 2 hrs - paths: - - build*/*Allen* - - build*/Sequence.json - - input - retry: 1 - -docker_image:build: - stage: build - only: - refs: - - master - tags: - - docker-image-build - script: "echo 'Building Allen dev docker image" - variables: - TO: $CI_REGISTRY_IMAGE:latest - allow_failure: true - -.run_physics_efficiency_job: &run_physics_efficiency_job_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: run - script: - - TOPLEVEL=${PWD} - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - declare -A DEVICE_MEMORY_MAP=${DEVICE_MEMORY} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=':' read -ra CI_RUNNER_DESCRIPTION_SPLIT <<< "${CI_RUNNER_DESCRIPTION}" - - IFS=${PREVIOUS_IFS} - - DEVICE_ID=${JOB_NAME_SPLIT[0]} - - TARGET=${JOB_NAME_SPLIT[1]} - - SEQUENCE=${JOB_NAME_SPLIT[2]} - - D_NUMBER=${CI_RUNNER_DESCRIPTION_SPLIT[1]} - - RUN_OPTIONS="-n 1000 -m 1000" - - export PATH=$PATH:/usr/local/cuda/bin - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - mkdir validation_output - - ln -s validation_output output # Needed by Root build - - cd build_${TARGET} - - ls - - export LD_LIBRARY_PATH=${PWD}:$LD_LIBRARY_PATH - - CUDA_VISIBLE_DEVICES=${D_NUMBER} numactl --cpunodebind=0 --membind=0 ./Allen -f /scratch/dcampora/allen_data/201907/bsphiphi_mag_down ${RUN_OPTIONS} 2>&1 | tee ../validation_output/bsphiphi_${DEVICE_ID}.txt - artifacts: - expire_in: 2 hrs - paths: - - validation_output/* - allow_failure: true - retry: 1 - -.run_physics_efficiency_job: &run_physics_efficiency_job_def_cpu - only: - refs: - - master - - schedules - - web - - merge_requests - stage: run - script: - - TOPLEVEL=${PWD} - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - declare -A DEVICE_MEMORY_MAP=${DEVICE_MEMORY} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=':' read -ra CI_RUNNER_DESCRIPTION_SPLIT <<< "${CI_RUNNER_DESCRIPTION}" - - IFS=${PREVIOUS_IFS} - - DEVICE_ID=${JOB_NAME_SPLIT[0]} - - TARGET=${JOB_NAME_SPLIT[1]} - - SEQUENCE=${JOB_NAME_SPLIT[2]} - - NUMA_NODE=${CI_RUNNER_DESCRIPTION_SPLIT[2]} - - RUN_OPTIONS="-n 1000 -m 1000" - - export PATH=$PATH:/usr/local/cuda/bin - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - mkdir validation_output - - ln -s validation_output output # Needed by Root build - - cd build_${TARGET} - - ls - - export LD_LIBRARY_PATH=${PWD}:$LD_LIBRARY_PATH - - numactl --cpunodebind=${NUMA_NODE} --membind=${NUMA_NODE} ./Allen -f /scratch/dcampora/allen_data/201907/bsphiphi_mag_down ${RUN_OPTIONS} 2>&1 | tee ../validation_output/bsphiphi_${DEVICE_ID}.txt - artifacts: - expire_in: 2 hrs - paths: - - validation_output/* - allow_failure: true - retry: 1 - -.run_throughput_job_no_profiling: &run_throughput_job_no_profiling_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: run - script: - - TOPLEVEL=${PWD} - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - declare -A DEVICE_MEMORY_MAP=${DEVICE_MEMORY} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=':' read -ra CI_RUNNER_DESCRIPTION_SPLIT <<< "${CI_RUNNER_DESCRIPTION}" - - IFS=${PREVIOUS_IFS} - - DEVICE_ID=${JOB_NAME_SPLIT[0]} - - TARGET=${JOB_NAME_SPLIT[1]} - - SEQUENCE=${JOB_NAME_SPLIT[2]} - - D_NUMBER=${CI_RUNNER_DESCRIPTION_SPLIT[1]} - - D_MEMORY=${DEVICE_MEMORY_MAP[${DEVICE_ID}]} - - RUN_OPTIONS="-n 500 -m 500 -r 1000 -t 16 -c 0" - - if [ "${D_MEMORY}" = "LOW" ]; then - - RUN_OPTIONS="-n 500 -m 500 -r 1000 -t 8 -c 0" - - fi - - export PATH=$PATH:/usr/local/cuda/bin - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - mkdir output_${DEVICE_ID} - - cd build_${TARGET} - - ls - - export LD_LIBRARY_PATH=${PWD}:$LD_LIBRARY_PATH - - CUDA_VISIBLE_DEVICES=${D_NUMBER} numactl --cpunodebind=0 --membind=0 ./Allen -f /scratch/dcampora/allen_data/201907/minbias_mag_down ${RUN_OPTIONS} 2>&1 | tee ../output_${DEVICE_ID}/output.txt - artifacts: - expire_in: 2 hrs - paths: - - output_* - allow_failure: true - retry: 1 - -.run_throughput_job_no_profiling_cpu: &run_throughput_job_no_profiling_def_cpu - only: - refs: - - master - - schedules - - web - - merge_requests - stage: run - script: - - TOPLEVEL=${PWD} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=':' read -ra CI_RUNNER_DESCRIPTION_SPLIT <<< "${CI_RUNNER_DESCRIPTION}" - - IFS=${PREVIOUS_IFS} - - DEVICE_ID=${JOB_NAME_SPLIT[0]} - - TARGET=${JOB_NAME_SPLIT[1]} - - SEQUENCE=${JOB_NAME_SPLIT[2]} - - TOTAL_THREADS=$(lscpu | egrep "^CPU\(s\):.*[0-9]+$" --color=none | awk '{ print $2; }') - - TOTAL_NUMA_NODES=$(lscpu | egrep "^NUMA node\(s\):.*[0-9]+$" --color=none | awk '{ print $3; }') - - NUMA_NODE=${CI_RUNNER_DESCRIPTION_SPLIT[2]} - - THREADS=$((${TOTAL_THREADS} / ${TOTAL_NUMA_NODES})) - - RUN_OPTIONS="-n 100 -m 100 -r 200 -t ${THREADS} -c 0" - - export PATH=$PATH:/usr/local/cuda/bin - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - mkdir output_${DEVICE_ID} - - cd build_${TARGET} - - ls - - export LD_LIBRARY_PATH=${PWD}:$LD_LIBRARY_PATH - - numactl --cpunodebind=${NUMA_NODE} --membind=${NUMA_NODE} ./Allen -f /scratch/dcampora/allen_data/201907/minbias_mag_down ${RUN_OPTIONS} 2>&1 | tee ../output_${DEVICE_ID}/output.txt - artifacts: - expire_in: 2 hrs - paths: - - output_* - allow_failure: true - retry: 1 - -.run_throughput_job: &run_throughput_job_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: run - script: - - TOPLEVEL=${PWD} - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - declare -A DEVICE_MEMORY_MAP=${DEVICE_MEMORY} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=':' read -ra CI_RUNNER_DESCRIPTION_SPLIT <<< "${CI_RUNNER_DESCRIPTION}" - - IFS=${PREVIOUS_IFS} - - DEVICE_ID=${JOB_NAME_SPLIT[0]} - - TARGET=${JOB_NAME_SPLIT[1]} - - SEQUENCE=${JOB_NAME_SPLIT[2]} - - D_NUMBER=${CI_RUNNER_DESCRIPTION_SPLIT[1]} - - D_MEMORY=${DEVICE_MEMORY_MAP[${DEVICE_ID}]} - - RUN_OPTIONS="-n 500 -m 500 -r 1000 -t 16 -c 0" - - if [ "${D_MEMORY}" = "LOW" ]; then - - RUN_OPTIONS="-n 500 -m 500 -r 1000 -t 8 -c 0" - - fi - - export PATH=$PATH:/usr/local/cuda/bin - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - mkdir output_${DEVICE_ID} - - cd build_${TARGET} - - ls - - export LD_LIBRARY_PATH=${PWD}:$LD_LIBRARY_PATH - - CUDA_VISIBLE_DEVICES=${D_NUMBER} ./Allen -f /scratch/dcampora/allen_data/201907/minbias_mag_down ${RUN_OPTIONS} 2>&1 | tee ../output_${DEVICE_ID}/output.txt - - NVPROF_TMPDIR=`mktemp -d` - - CUDA_VISIBLE_DEVICES=${D_NUMBER} TMPDIR=${NVPROF_TMPDIR} numactl --cpunodebind=0 --membind=0 nvprof ./Allen -f /scratch/dcampora/allen_data/201907/minbias_mag_down ${RUN_OPTIONS} 2>&1 | tee ../output_${DEVICE_ID}/profiler_output.txt - - rm -rf ${NVPROF_TMPDIR} - - python3 ${TOPLEVEL}/checker/plotting/extract_algo_breakdown.py -d ${TOPLEVEL} - artifacts: - expire_in: 2 hrs - paths: - - output_* - allow_failure: true - retry: 1 - -.throughput_cli_plot_job: &publish_algo_breakdown_plot_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: publish - script: - - declare -A DEVICE_NUMBERS_MAP=${DEVICE_NUMBERS} - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=${PREVIOUS_IFS} - - DEVICE_ID=${JOB_NAME_SPLIT[0]} - - SEQUENCE=${JOB_NAME_SPLIT[1]} - - export PATH=$PATH:/usr/local/cuda/bin - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - python3 checker/plotting/csv_plotter.py -t "Algorithm Breakdown of sequence ${SEQUENCE}, branch ${CI_COMMIT_REF_NAME}" -u "%" -x 30 -m ${MATTERMOST_KEY} output_${DEVICE_ID}/algo_breakdown.csv - -test_physics_efficiency: - only: - refs: - - master - - schedules - - web - - merge_requests - stage: test - script: - - TOPLEVEL=${PWD} - - ls validation_output - - ls ${TOPLEVEL}/test/reference - - cd validation_output - - for i in $( ls ); do echo "Checking ${i}"; tail -n97 ${i} | head -n94 > efficiency_${i}; diff -u ${TOPLEVEL}/test/reference/${i} efficiency_${i} | tee ${i}_diff || true; done - - cat *_diff > alldiffs - - if [ -s alldiffs ]; then echo "Differences were found against reference files."; exit 1; else echo "No differences found against reference files."; exit 0; fi - dependencies: - - geforcertx2080ti:CUDA:hlt1_pp_default:run_physics_efficiency - - x862630v4:CPU:hlt1_pp_default:run_physics_efficiency - tags: - - cvmfs - allow_failure: true - -# run_built_tests: -# only: -# refs: -# - master -# - schedules -# - web -# - merge_requests -# stage: test -# script: -# - cd build_CUDA -# - ctest -V -# dependencies: -# - CUDA:hlt1_pp_default:Debug::build -# allow_failure: true - -.publish_throughput_job: &publish_throughput_job_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: publish - script: - - PREVIOUS_IFS=${IFS} - - IFS=':' read -ra JOB_NAME_SPLIT <<< "${CI_JOB_NAME}" - - IFS=${PREVIOUS_IFS} - - SEQUENCE=${JOB_NAME_SPLIT[1]} - - BREAKDOWN_DEVICE_ID=${JOB_NAME_SPLIT[2]} - - cat output_*/output.txt | grep --color=none "device" | sed 's/.*:\ [0-9]*\,\ //' > devices.txt - - cat output_*/output.txt | grep --color=none "events/s" | awk '{ print $1; }' > throughputs.txt - - cat devices.txt - - cat throughputs.txt - - paste -d, devices.txt throughputs.txt > devices_throughputs.csv - - cat devices_throughputs.csv - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - python3 checker/plotting/post_combined_message.py -l "Throughput of [sequence ${SEQUENCE}, branch ${CI_COMMIT_REF_NAME}](https://gitlab.cern.ch/lhcb/Allen/pipelines/${CI_PIPELINE_ID})" -m ${MATTERMOST_KEY} -t devices_throughputs.csv -b output_${BREAKDOWN_DEVICE_ID}/algo_breakdown.csv - - python3 checker/plotting/post_telegraf.py -d . -s ${SEQUENCE} -b ${CI_COMMIT_REF_NAME} - -.throughput_speedup_job: &publish_speedup_job_def - only: - refs: - - master - - schedules - - web - - merge_requests - stage: publish - script: - - cat output_*/output.txt | grep --color=none "device" | sed 's/.*:\ [0-9]*\,\ //' > devices.txt - - cat output_*/output.txt | grep --color=none "events/s" | awk '{ print $1; }' > throughputs.txt - - cat devices.txt - - cat throughputs.txt - - paste -d, devices.txt throughputs.txt > devices_throughputs.csv - - cat devices_throughputs.csv - - source /cvmfs/sft.cern.ch/lcg/views/setupViews.sh LCG_97apython3 x86_64-centos7-gcc9-opt - - python3 checker/plotting/csv_plotter.py -n -t "Speedup across GPUs, branch _${CI_COMMIT_REF_NAME}_" -u "x" -x 30 -m ${MATTERMOST_KEY} devices_throughputs.csv - -# ===== -# Build -# ===== - -CUDA:hlt1_pp_default:RelWithDebInfo::build: - <<: *build_clang_job_def - tags: - - cvmfs - -CUDA:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build: - <<: *build_clang_job_def - tags: - - cvmfs - -CPU:hlt1_pp_default:RelWithDebInfo::build: - <<: *build_clang_job_def - tags: - - cvmfs - -CPU:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build: - <<: *build_clang_job_def - tags: - - cvmfs - -# Builds with warnings treated as errors -# gcc: - -CPU:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build_warnings_as_errors: - <<: *build_gcc_job_def_warnings_as_errors - tags: - - cvmfs - -CUDA:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build_warnings_as_errors: - <<: *build_gcc_job_def_warnings_as_errors - tags: - - cvmfs - -# clang: - -CPU:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build_warnings_as_errors: - <<: *build_clang_job_def_warnings_as_errors - tags: - - cvmfs - -CUDA:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build_warnings_as_errors: - <<: *build_clang_job_def_warnings_as_errors - tags: - - cvmfs - -# Builds only on master of additional options - -CUDA:hlt1_pp_no_gec:RelWithDebInfo::build: - <<: *build_clang_job_def - tags: - - cvmfs - only: - refs: - - master - -CPU:hlt1_pp_no_gec:RelWithDebInfo::build: - <<: *build_clang_job_def - tags: - - cvmfs - only: - refs: - - master - -# === -# Run -# === - -# Throughput runs - -geforcertx2080ti:CUDA:hlt1_pp_default:run_throughput: - <<: *run_throughput_job_def - tags: - - geforcertx2080ti - dependencies: - - CUDA:hlt1_pp_default:RelWithDebInfo::build - -quadrortx6000:CUDA:hlt1_pp_default:run_throughput: - <<: *run_throughput_job_no_profiling_def - tags: - - quadrortx6000 - dependencies: - - CUDA:hlt1_pp_default:RelWithDebInfo::build - -teslav100:CUDA:hlt1_pp_default:run_throughput: - <<: *run_throughput_job_no_profiling_def - tags: - - teslav100 - dependencies: - - CUDA:hlt1_pp_default:RelWithDebInfo::build - -# teslat4:CUDA:hlt1_pp_default:run_throughput: -# <<: *run_throughput_job_no_profiling_def -# tags: -# - teslat4 -# dependencies: -# - CUDA:hlt1_pp_default:RelWithDebInfo::build - -x862630v4:CPU:hlt1_pp_default:run_throughput_cpu: - <<: *run_throughput_job_no_profiling_def_cpu - tags: - - x862630v4 - dependencies: - - CPU:hlt1_pp_default:RelWithDebInfo::build - -epyc7502:CPU:hlt1_pp_default:run_throughput_cpu: - <<: *run_throughput_job_no_profiling_def_cpu - tags: - - epyc7502 - dependencies: - - CPU:hlt1_pp_default:RelWithDebInfo::build - -# Physics efficiency runs - -geforcertx2080ti:CUDA:hlt1_pp_default:run_physics_efficiency: - <<: *run_physics_efficiency_job_def - tags: - - geforcertx2080ti - dependencies: - - CUDA:hlt1_pp_default:RelWithDebInfo::build - -x862630v4:CPU:hlt1_pp_default:run_physics_efficiency: - <<: *run_physics_efficiency_job_def_cpu - tags: - - x862630v4 - dependencies: - - CPU:hlt1_pp_default:RelWithDebInfo::build - -# Test runs - -# geforcertx2080ti:CUDA:hlt1_pp_default:run_physics_efficiency_debug: -# <<: *run_physics_efficiency_job_def -# tags: -# - geforcertx2080ti -# dependencies: -# - CUDA:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build - -# x86:CPU:hlt1_pp_default:run_physics_efficiency_debug: -# <<: *run_physics_efficiency_job_def_cpu -# tags: -# - x86 -# dependencies: -# - CPU:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build - -# x86:CPU:hlt1_pp_default:run_physics_efficiency_debug_root: -# <<: *run_physics_efficiency_job_def_cpu -# tags: -# - x86 -# dependencies: -# - CPU:hlt1_pp_default:Debug:-DUSE_ROOT=ON -DBUILD_TESTS=ON:build - -# Additional throughput runs, only on master - -geforcertx2080ti:CUDA:hlt1_pp_no_gec:run_throughput: - <<: *run_throughput_job_def - tags: - - geforcertx2080ti - dependencies: - - CUDA:hlt1_pp_no_gec:RelWithDebInfo::build - only: - refs: - - master - -quadrortx6000:CUDA:hlt1_pp_no_gec:run_throughput: - <<: *run_throughput_job_no_profiling_def - tags: - - quadrortx6000 - dependencies: - - CUDA:hlt1_pp_no_gec:RelWithDebInfo::build - only: - refs: - - master - -teslav100:CUDA:hlt1_pp_no_gec:run_throughput: - <<: *run_throughput_job_no_profiling_def - tags: - - teslav100 - dependencies: - - CUDA:hlt1_pp_no_gec:RelWithDebInfo::build - only: - refs: - - master - -# teslat4:CUDA:hlt1_pp_no_gec:run_throughput: -# <<: *run_throughput_job_no_profiling_def -# tags: -# - teslat4 -# dependencies: -# - CUDA:hlt1_pp_no_gec:RelWithDebInfo::build -# only: -# refs: -# - master - -x862630v4:CPU:hlt1_pp_no_gec:run_throughput_cpu: - <<: *run_throughput_job_no_profiling_def_cpu - tags: - - x862630v4 - dependencies: - - CPU:hlt1_pp_no_gec:RelWithDebInfo::build - only: - refs: - - master - -epyc7502:CPU:hlt1_pp_no_gec:run_throughput_cpu: - <<: *run_throughput_job_no_profiling_def_cpu - tags: - - epyc7502 - dependencies: - - CPU:hlt1_pp_no_gec:RelWithDebInfo::build - only: - refs: - - master - -# ======= -# Publish -# ======= - -throughput:hlt1_pp_default:geforcertx2080ti:publish_throughput: - <<: *publish_throughput_job_def - tags: - - cvmfs - dependencies: - - geforcertx2080ti:CUDA:hlt1_pp_default:run_throughput - - quadrortx6000:CUDA:hlt1_pp_default:run_throughput - # - teslat4:CUDA:hlt1_pp_default:run_throughput - - teslav100:CUDA:hlt1_pp_default:run_throughput - - x862630v4:CPU:hlt1_pp_default:run_throughput_cpu - - epyc7502:CPU:hlt1_pp_default:run_throughput_cpu - -# Publish additional throughput tests, only in master - -throughput:hlt1_pp_no_gec:geforcertx2080ti:publish_throughput: - <<: *publish_throughput_job_def - tags: - - x86 - dependencies: - - geforcertx2080ti:CUDA:hlt1_pp_no_gec:run_throughput - - quadrortx6000:CUDA:hlt1_pp_no_gec:run_throughput - # - teslat4:CUDA:hlt1_pp_no_gec:run_throughput - - teslav100:CUDA:hlt1_pp_no_gec:run_throughput - - x862630v4:CPU:hlt1_pp_no_gec:run_throughput_cpu - - epyc7502:CPU:hlt1_pp_no_gec:run_throughput_cpu - only: - refs: - - master diff --git a/Dumpers/BinaryDumpers/options/dump.py b/Dumpers/BinaryDumpers/options/dump.py index 66c628c27d12cd07c0b7358e584d234eca7ab74f..7a11b61865503fe82535ea4d93daa02f814f9974 100644 --- a/Dumpers/BinaryDumpers/options/dump.py +++ b/Dumpers/BinaryDumpers/options/dump.py @@ -63,7 +63,7 @@ ApplicationMgr().ExtSvc += [ ] # Dump raw banks and UT, FT and muon hits -transpose_banks = TransposeRawBanks(BankTypes=["VP", "UT", "FTCluster", +transpose_banks = TransposeRawBanks(BankTypes=["VP", "VPRetinaCluster", "UT", "FTCluster", "Muon", "ODIN"]) dump_banks = DumpRawBanks() dump_ut = DumpUTHits() diff --git a/Dumpers/BinaryDumpers/options/dump_banks.py b/Dumpers/BinaryDumpers/options/dump_banks.py index 1106de77dd950b0b4b40a45f498c60516856e990..6a055d766b5be0fd65afd316bd7b9bd24ba9aee3 100644 --- a/Dumpers/BinaryDumpers/options/dump_banks.py +++ b/Dumpers/BinaryDumpers/options/dump_banks.py @@ -60,7 +60,7 @@ ApplicationMgr().ExtSvc += [DumpUTGeometry()] # Dump raw banks and UT, FT and muon hits transpose_banks = TransposeRawBanks( - BankTypes=["VP", "UT", "FTCluster", "Muon", "ODIN"], + BankTypes=["VP", "VPRetinaCluster", "UT", "FTCluster", "Muon", "ODIN"], RawEventLocations=["DAQ/RawEvent"]) dump_banks = DumpRawBanks() dump_seq = GaudiSequencer("DumpSeq") diff --git a/Dumpers/BinaryDumpers/options/dump_banks_and_MC_info.py b/Dumpers/BinaryDumpers/options/dump_banks_and_MC_info.py index 38def5917705b89031e18ab98f19059b753582ec..4d139f1604229aa62aa7dbad4fb6bb97423e9b75 100644 --- a/Dumpers/BinaryDumpers/options/dump_banks_and_MC_info.py +++ b/Dumpers/BinaryDumpers/options/dump_banks_and_MC_info.py @@ -92,7 +92,8 @@ def AddDumpers(): dump_mc.MCOutputDirectory = output_file + "/MC_info/tracks" dump_pvmc = PVDumper("DumpPVMCInfo") dump_pvmc.OutputDirectory = output_file + "/MC_info/PVs" - dump_banks = DumpRawBanks(BankTypes=["VP", "UT", "FTCluster", "Muon"]) + dump_banks = DumpRawBanks( + BankTypes=["VP", "VPRetinaCluster", "UT", "FTCluster", "Muon"]) dump_banks.OutputDirectory = output_file + "/banks" dump_muon_coords = DumpMuonCoords() dump_muon_coords.OutputDirectory = output_file + "/muon_coords" diff --git a/Rec/Allen/options/run_allen_in_gaudi.py b/Rec/Allen/options/run_allen_in_gaudi.py index bb670a34ed9004a46a03db91f124547e1c776dd8..c9b778282510b159ba3ff3e223a1cd2a91624fca 100644 --- a/Rec/Allen/options/run_allen_in_gaudi.py +++ b/Rec/Allen/options/run_allen_in_gaudi.py @@ -156,7 +156,7 @@ NTupleSvc().Output = ["FILE1 DATAFILE='velo_states.root' TYP='ROOT' OPT='NEW'"] # Save raw banks in Allen format on the TES outputdirectory = "dump/" transpose_banks = TransposeRawBanks( - BankTypes=["VP", "UT", "FTCluster", "Muon", "ODIN"], + BankTypes=["VP", "VPRetinaCluster", "UT", "FTCluster", "Muon", "ODIN"], RawEventLocations=["DAQ/RawEvent"]) dump_seq = GaudiSequencer("RecoAllenPrepareSeq") dump_seq.Members += [transpose_banks] diff --git a/Rec/Allen/src/RunAllen.cpp b/Rec/Allen/src/RunAllen.cpp index 3723470f145b8af3f4ec7c9d53a48fb08bfbf39f..9d5c0a6517c27cef73693d2ada0aeb64dae2c513 100644 --- a/Rec/Allen/src/RunAllen.cpp +++ b/Rec/Allen/src/RunAllen.cpp @@ -115,9 +115,13 @@ StatusCode RunAllen::initialize() const size_t number_of_slices = 1; const size_t events_per_slice = 1; const size_t n_events = 1; - m_tes_input_provider.reset( - new TESProvider( - number_of_slices, events_per_slice, n_events)); + m_tes_input_provider.reset(new TESProvider< + BankTypes::VP, + BankTypes::VPRetinaCluster, + BankTypes::UT, + BankTypes::FT, + BankTypes::MUON, + BankTypes::ODIN>(number_of_slices, events_per_slice, n_events)); // Get HLT1 selection names from configuration and initialize rate counters m_line_names = configuration_reader.params()["configured_lines"]; diff --git a/Rec/Allen/src/RunAllen.h b/Rec/Allen/src/RunAllen.h index e75198613c5589a5c5dd4a9e862cfcbc224e32b0..d60ac017c3260fe77513a7d5b0a56af53ee2e631 100644 --- a/Rec/Allen/src/RunAllen.h +++ b/Rec/Allen/src/RunAllen.h @@ -62,6 +62,7 @@ private: Constants m_constants; std::set m_bankTypes = {LHCb::RawBank::ODIN, LHCb::RawBank::VP, + LHCb::RawBank::VPRetinaCluster, LHCb::RawBank::UT, LHCb::RawBank::FTCluster, LHCb::RawBank::Muon}; @@ -74,7 +75,13 @@ private: std::unique_ptr m_stream_wrapper; std::unique_ptr m_host_buffers_manager; - std::unique_ptr> + std::unique_ptr> m_tes_input_provider; Gaudi::Property m_updaterName {this, "UpdaterName", "AllenUpdater"}; diff --git a/configuration/sequences/definitions/VeloSequence.py b/configuration/sequences/definitions/VeloSequence.py index 615c0f629aca34017a9ed7ebc4053b38e2aee906..cd97bf1d7bf771be9342e4611bbe1fa88f4a58e7 100644 --- a/configuration/sequences/definitions/VeloSequence.py +++ b/configuration/sequences/definitions/VeloSequence.py @@ -1,7 +1,7 @@ from definitions.algorithms import * -def VeloSequence(doGEC=True): +def VeloSequence(doGEC=True, retina_decoding=False): host_ut_banks = host_data_provider_t(name="host_ut_banks", bank_type="UT") host_scifi_banks = host_data_provider_t( @@ -23,162 +23,303 @@ def VeloSequence(doGEC=True): host_scifi_raw_banks_t=host_scifi_banks.host_raw_banks_t(), host_scifi_raw_offsets_t=host_scifi_banks.host_raw_offsets_t()) - velo_banks = data_provider_t(name="velo_banks", bank_type="VP") - - velo_calculate_number_of_candidates = velo_calculate_number_of_candidates_t( - name="velo_calculate_number_of_candidates", - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - dev_event_list_t=initialize_lists.dev_event_list_t(), - dev_velo_raw_input_t=velo_banks.dev_raw_banks_t(), - dev_velo_raw_input_offsets_t=velo_banks.dev_raw_offsets_t()) - - prefix_sum_offsets_velo_candidates = host_prefix_sum_t( - name="prefix_sum_offsets_velo_candidates", - dev_input_buffer_t=velo_calculate_number_of_candidates. - dev_number_of_candidates_t()) - - velo_estimate_input_size = velo_estimate_input_size_t( - name="velo_estimate_input_size", - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - host_number_of_cluster_candidates_t=prefix_sum_offsets_velo_candidates. - host_total_sum_holder_t(), - dev_event_list_t=initialize_lists.dev_event_list_t(), - dev_candidates_offsets_t=prefix_sum_offsets_velo_candidates. - dev_output_buffer_t(), - dev_velo_raw_input_t=velo_banks.dev_raw_banks_t(), - dev_velo_raw_input_offsets_t=velo_banks.dev_raw_offsets_t()) - - prefix_sum_offsets_estimated_input_size = host_prefix_sum_t( - name="prefix_sum_offsets_estimated_input_size", - dev_input_buffer_t=velo_estimate_input_size. - dev_estimated_input_size_t()) - - velo_masked_clustering = velo_masked_clustering_t( - name="velo_masked_clustering", - host_total_number_of_velo_clusters_t= - prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - dev_velo_raw_input_t=velo_banks.dev_raw_banks_t(), - dev_velo_raw_input_offsets_t=velo_banks.dev_raw_offsets_t(), - dev_offsets_estimated_input_size_t= - prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), - dev_module_candidate_num_t=velo_estimate_input_size. - dev_module_candidate_num_t(), - dev_cluster_candidates_t=velo_estimate_input_size. - dev_cluster_candidates_t(), - dev_event_list_t=initialize_lists.dev_event_list_t(), - dev_candidates_offsets_t=prefix_sum_offsets_velo_candidates. - dev_output_buffer_t()) - - velo_calculate_phi_and_sort = velo_calculate_phi_and_sort_t( - name="velo_calculate_phi_and_sort", - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - host_total_number_of_velo_clusters_t= - prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), - dev_offsets_estimated_input_size_t= - prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), - dev_module_cluster_num_t=velo_masked_clustering. - dev_module_cluster_num_t(), - dev_velo_cluster_container_t=velo_masked_clustering. - dev_velo_cluster_container_t()) - - velo_search_by_triplet = velo_search_by_triplet_t( - name="velo_search_by_triplet", - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - host_total_number_of_velo_clusters_t= - prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), - dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. - dev_sorted_velo_cluster_container_t(), - dev_offsets_estimated_input_size_t= - prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), - dev_module_cluster_num_t=velo_masked_clustering. - dev_module_cluster_num_t(), - dev_hit_phi_t=velo_calculate_phi_and_sort.dev_hit_phi_t()) - - prefix_sum_offsets_velo_tracks = host_prefix_sum_t( - name="prefix_sum_offsets_velo_tracks", - dev_input_buffer_t=velo_search_by_triplet. - dev_number_of_velo_tracks_t()) - - velo_three_hit_tracks_filter = velo_three_hit_tracks_filter_t( - name="velo_three_hit_tracks_filter", - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. - dev_sorted_velo_cluster_container_t(), - dev_offsets_estimated_input_size_t= - prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), - dev_atomics_velo_t=velo_search_by_triplet.dev_atomics_velo_t(), - dev_hit_used_t=velo_search_by_triplet.dev_hit_used_t(), - dev_three_hit_tracks_input_t=velo_search_by_triplet. - dev_three_hit_tracks_t()) - - prefix_sum_offsets_number_of_three_hit_tracks_filtered = host_prefix_sum_t( - name="prefix_sum_offsets_number_of_three_hit_tracks_filtered", - dev_input_buffer_t=velo_three_hit_tracks_filter. - dev_number_of_three_hit_tracks_output_t()) - - velo_copy_track_hit_number = velo_copy_track_hit_number_t( - name="velo_copy_track_hit_number", - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - host_number_of_velo_tracks_at_least_four_hits_t= - prefix_sum_offsets_velo_tracks.host_total_sum_holder_t(), - host_number_of_three_hit_tracks_filtered_t= - prefix_sum_offsets_number_of_three_hit_tracks_filtered. - host_total_sum_holder_t(), - dev_tracks_t=velo_search_by_triplet.dev_tracks_t(), - dev_offsets_velo_tracks_t=prefix_sum_offsets_velo_tracks. - dev_output_buffer_t(), - dev_offsets_number_of_three_hit_tracks_filtered_t= - prefix_sum_offsets_number_of_three_hit_tracks_filtered. - dev_output_buffer_t()) - - prefix_sum_offsets_velo_track_hit_number = host_prefix_sum_t( - name="prefix_sum_offsets_velo_track_hit_number", - dev_input_buffer_t=velo_copy_track_hit_number. - dev_velo_track_hit_number_t()) - - velo_consolidate_tracks = velo_consolidate_tracks_t( - name="velo_consolidate_tracks", - host_accumulated_number_of_hits_in_velo_tracks_t= - prefix_sum_offsets_velo_track_hit_number.host_total_sum_holder_t(), - host_number_of_reconstructed_velo_tracks_t=velo_copy_track_hit_number. - host_number_of_reconstructed_velo_tracks_t(), - host_number_of_three_hit_tracks_filtered_t= - prefix_sum_offsets_number_of_three_hit_tracks_filtered. - host_total_sum_holder_t(), - host_number_of_selected_events_t=initialize_lists. - host_number_of_selected_events_t(), - dev_offsets_all_velo_tracks_t=velo_copy_track_hit_number. - dev_offsets_all_velo_tracks_t(), - dev_tracks_t=velo_search_by_triplet.dev_tracks_t(), - dev_offsets_velo_track_hit_number_t= - prefix_sum_offsets_velo_track_hit_number.dev_output_buffer_t(), - dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. - dev_sorted_velo_cluster_container_t(), - dev_offsets_estimated_input_size_t= - prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), - dev_three_hit_tracks_output_t=velo_three_hit_tracks_filter. - dev_three_hit_tracks_output_t(), - dev_offsets_number_of_three_hit_tracks_filtered_t= - prefix_sum_offsets_number_of_three_hit_tracks_filtered. - dev_output_buffer_t()) - - velo_sequence = Sequence( - host_ut_banks, host_scifi_banks, initialize_lists, velo_banks, - velo_calculate_number_of_candidates, - prefix_sum_offsets_velo_candidates, velo_estimate_input_size, - prefix_sum_offsets_estimated_input_size, velo_masked_clustering, - velo_calculate_phi_and_sort, velo_search_by_triplet, - prefix_sum_offsets_velo_tracks, velo_three_hit_tracks_filter, - prefix_sum_offsets_number_of_three_hit_tracks_filtered, - velo_copy_track_hit_number, prefix_sum_offsets_velo_track_hit_number, - velo_consolidate_tracks) + if retina_decoding: + velo_retina_banks = data_provider_t( + name="velo_retina_banks", bank_type="VPRetinaCluster") + + calculate_number_of_retinaclusters_each_sensor = calculate_number_of_retinaclusters_each_sensor_t( + name="calculate_number_of_retinaclusters_each_sensor", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_event_list_t=initialize_lists.dev_event_list_t(), + dev_velo_retina_raw_input_t=velo_retina_banks.dev_raw_banks_t(), + dev_velo_retina_raw_input_offsets_t=velo_retina_banks. + dev_raw_offsets_t()) + + prefix_sum_offsets_estimated_input_size = host_prefix_sum_t( + name="prefix_sum_offsets_estimated_input_size", + dev_input_buffer_t= + calculate_number_of_retinaclusters_each_sensor. + dev_each_sensor_size_t()) + + decode_retinaclusters = decode_retinaclusters_t( + name="decode_retinaclusters", + host_total_number_of_velo_clusters_t= + prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_velo_retina_raw_input_t=velo_retina_banks.dev_raw_banks_t(), + dev_velo_retina_raw_input_offsets_t=velo_retina_banks. + dev_raw_offsets_t(), + dev_offsets_each_sensor_size_t= + prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), + dev_event_list_t=initialize_lists.dev_event_list_t()) + + velo_calculate_phi_and_sort = velo_calculate_phi_and_sort_t( + name="velo_calculate_phi_and_sort", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + host_total_number_of_velo_clusters_t= + prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), + dev_offsets_estimated_input_size_t=decode_retinaclusters. + dev_offsets_module_pair_cluster_t(), + dev_module_cluster_num_t=decode_retinaclusters. + dev_module_cluster_num_t(), + dev_velo_cluster_container_t=decode_retinaclusters. + dev_velo_cluster_container_t()) + + velo_search_by_triplet = velo_search_by_triplet_t( + name="velo_search_by_triplet", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + host_total_number_of_velo_clusters_t= + prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), + dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. + dev_sorted_velo_cluster_container_t(), + dev_offsets_estimated_input_size_t=decode_retinaclusters. + dev_offsets_module_pair_cluster_t(), + dev_module_cluster_num_t=decode_retinaclusters. + dev_module_cluster_num_t(), + dev_hit_phi_t=velo_calculate_phi_and_sort.dev_hit_phi_t()) + + prefix_sum_offsets_velo_tracks = host_prefix_sum_t( + name="prefix_sum_offsets_velo_tracks", + dev_input_buffer_t=velo_search_by_triplet. + dev_number_of_velo_tracks_t()) + + velo_three_hit_tracks_filter = velo_three_hit_tracks_filter_t( + name="velo_three_hit_tracks_filter", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. + dev_sorted_velo_cluster_container_t(), + dev_offsets_estimated_input_size_t=decode_retinaclusters. + dev_offsets_module_pair_cluster_t(), + dev_atomics_velo_t=velo_search_by_triplet.dev_atomics_velo_t(), + dev_hit_used_t=velo_search_by_triplet.dev_hit_used_t(), + dev_three_hit_tracks_input_t=velo_search_by_triplet. + dev_three_hit_tracks_t()) + + prefix_sum_offsets_number_of_three_hit_tracks_filtered = host_prefix_sum_t( + name="prefix_sum_offsets_number_of_three_hit_tracks_filtered", + dev_input_buffer_t=velo_three_hit_tracks_filter. + dev_number_of_three_hit_tracks_output_t()) + + velo_copy_track_hit_number = velo_copy_track_hit_number_t( + name="velo_copy_track_hit_number", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + host_number_of_velo_tracks_at_least_four_hits_t= + prefix_sum_offsets_velo_tracks.host_total_sum_holder_t(), + host_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + host_total_sum_holder_t(), + dev_tracks_t=velo_search_by_triplet.dev_tracks_t(), + dev_offsets_velo_tracks_t=prefix_sum_offsets_velo_tracks. + dev_output_buffer_t(), + dev_offsets_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + dev_output_buffer_t()) + + prefix_sum_offsets_velo_track_hit_number = host_prefix_sum_t( + name="prefix_sum_offsets_velo_track_hit_number", + dev_input_buffer_t=velo_copy_track_hit_number. + dev_velo_track_hit_number_t()) + + velo_consolidate_tracks = velo_consolidate_tracks_t( + name="velo_consolidate_tracks", + host_accumulated_number_of_hits_in_velo_tracks_t= + prefix_sum_offsets_velo_track_hit_number.host_total_sum_holder_t(), + host_number_of_reconstructed_velo_tracks_t=velo_copy_track_hit_number. + host_number_of_reconstructed_velo_tracks_t(), + host_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + host_total_sum_holder_t(), + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_offsets_all_velo_tracks_t=velo_copy_track_hit_number. + dev_offsets_all_velo_tracks_t(), + dev_tracks_t=velo_search_by_triplet.dev_tracks_t(), + dev_offsets_velo_track_hit_number_t= + prefix_sum_offsets_velo_track_hit_number.dev_output_buffer_t(), + dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. + dev_sorted_velo_cluster_container_t(), + dev_offsets_estimated_input_size_t=decode_retinaclusters. + dev_offsets_module_pair_cluster_t(), + dev_three_hit_tracks_output_t=velo_three_hit_tracks_filter. + dev_three_hit_tracks_output_t(), + dev_offsets_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + dev_output_buffer_t()) + else: + velo_banks = data_provider_t(name="velo_banks", bank_type="VP") + + velo_calculate_number_of_candidates = velo_calculate_number_of_candidates_t( + name="velo_calculate_number_of_candidates", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_event_list_t=initialize_lists.dev_event_list_t(), + dev_velo_raw_input_t=velo_banks.dev_raw_banks_t(), + dev_velo_raw_input_offsets_t=velo_banks.dev_raw_offsets_t()) + + prefix_sum_offsets_velo_candidates = host_prefix_sum_t( + name="prefix_sum_offsets_velo_candidates", + dev_input_buffer_t=velo_calculate_number_of_candidates. + dev_number_of_candidates_t()) + + velo_estimate_input_size = velo_estimate_input_size_t( + name="velo_estimate_input_size", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + host_number_of_cluster_candidates_t= + prefix_sum_offsets_velo_candidates.host_total_sum_holder_t(), + dev_event_list_t=initialize_lists.dev_event_list_t(), + dev_candidates_offsets_t=prefix_sum_offsets_velo_candidates. + dev_output_buffer_t(), + dev_velo_raw_input_t=velo_banks.dev_raw_banks_t(), + dev_velo_raw_input_offsets_t=velo_banks.dev_raw_offsets_t()) + + prefix_sum_offsets_estimated_input_size = host_prefix_sum_t( + name="prefix_sum_offsets_estimated_input_size", + dev_input_buffer_t=velo_estimate_input_size. + dev_estimated_input_size_t()) + + velo_masked_clustering = velo_masked_clustering_t( + name="velo_masked_clustering", + host_total_number_of_velo_clusters_t= + prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_velo_raw_input_t=velo_banks.dev_raw_banks_t(), + dev_velo_raw_input_offsets_t=velo_banks.dev_raw_offsets_t(), + dev_offsets_estimated_input_size_t= + prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), + dev_module_candidate_num_t=velo_estimate_input_size. + dev_module_candidate_num_t(), + dev_cluster_candidates_t=velo_estimate_input_size. + dev_cluster_candidates_t(), + dev_event_list_t=initialize_lists.dev_event_list_t(), + dev_candidates_offsets_t=prefix_sum_offsets_velo_candidates. + dev_output_buffer_t()) + + velo_calculate_phi_and_sort = velo_calculate_phi_and_sort_t( + name="velo_calculate_phi_and_sort", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + host_total_number_of_velo_clusters_t= + prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), + dev_offsets_estimated_input_size_t= + prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), + dev_module_cluster_num_t=velo_masked_clustering. + dev_module_cluster_num_t(), + dev_velo_cluster_container_t=velo_masked_clustering. + dev_velo_cluster_container_t()) + + velo_search_by_triplet = velo_search_by_triplet_t( + name="velo_search_by_triplet", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + host_total_number_of_velo_clusters_t= + prefix_sum_offsets_estimated_input_size.host_total_sum_holder_t(), + dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. + dev_sorted_velo_cluster_container_t(), + dev_offsets_estimated_input_size_t= + prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), + dev_module_cluster_num_t=velo_masked_clustering. + dev_module_cluster_num_t(), + dev_hit_phi_t=velo_calculate_phi_and_sort.dev_hit_phi_t()) + + prefix_sum_offsets_velo_tracks = host_prefix_sum_t( + name="prefix_sum_offsets_velo_tracks", + dev_input_buffer_t=velo_search_by_triplet. + dev_number_of_velo_tracks_t()) + + velo_three_hit_tracks_filter = velo_three_hit_tracks_filter_t( + name="velo_three_hit_tracks_filter", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. + dev_sorted_velo_cluster_container_t(), + dev_offsets_estimated_input_size_t= + prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), + dev_atomics_velo_t=velo_search_by_triplet.dev_atomics_velo_t(), + dev_hit_used_t=velo_search_by_triplet.dev_hit_used_t(), + dev_three_hit_tracks_input_t=velo_search_by_triplet. + dev_three_hit_tracks_t()) + + prefix_sum_offsets_number_of_three_hit_tracks_filtered = host_prefix_sum_t( + name="prefix_sum_offsets_number_of_three_hit_tracks_filtered", + dev_input_buffer_t=velo_three_hit_tracks_filter. + dev_number_of_three_hit_tracks_output_t()) + + velo_copy_track_hit_number = velo_copy_track_hit_number_t( + name="velo_copy_track_hit_number", + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + host_number_of_velo_tracks_at_least_four_hits_t= + prefix_sum_offsets_velo_tracks.host_total_sum_holder_t(), + host_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + host_total_sum_holder_t(), + dev_tracks_t=velo_search_by_triplet.dev_tracks_t(), + dev_offsets_velo_tracks_t=prefix_sum_offsets_velo_tracks. + dev_output_buffer_t(), + dev_offsets_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + dev_output_buffer_t()) + + prefix_sum_offsets_velo_track_hit_number = host_prefix_sum_t( + name="prefix_sum_offsets_velo_track_hit_number", + dev_input_buffer_t=velo_copy_track_hit_number. + dev_velo_track_hit_number_t()) + + velo_consolidate_tracks = velo_consolidate_tracks_t( + name="velo_consolidate_tracks", + host_accumulated_number_of_hits_in_velo_tracks_t= + prefix_sum_offsets_velo_track_hit_number.host_total_sum_holder_t(), + host_number_of_reconstructed_velo_tracks_t=velo_copy_track_hit_number. + host_number_of_reconstructed_velo_tracks_t(), + host_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + host_total_sum_holder_t(), + host_number_of_selected_events_t=initialize_lists. + host_number_of_selected_events_t(), + dev_offsets_all_velo_tracks_t=velo_copy_track_hit_number. + dev_offsets_all_velo_tracks_t(), + dev_tracks_t=velo_search_by_triplet.dev_tracks_t(), + dev_offsets_velo_track_hit_number_t= + prefix_sum_offsets_velo_track_hit_number.dev_output_buffer_t(), + dev_sorted_velo_cluster_container_t=velo_calculate_phi_and_sort. + dev_sorted_velo_cluster_container_t(), + dev_offsets_estimated_input_size_t= + prefix_sum_offsets_estimated_input_size.dev_output_buffer_t(), + dev_three_hit_tracks_output_t=velo_three_hit_tracks_filter. + dev_three_hit_tracks_output_t(), + dev_offsets_number_of_three_hit_tracks_filtered_t= + prefix_sum_offsets_number_of_three_hit_tracks_filtered. + dev_output_buffer_t()) + + if retina_decoding: + velo_sequence = Sequence( + host_ut_banks, host_scifi_banks, initialize_lists, + velo_retina_banks, + calculate_number_of_retinaclusters_each_sensor, + prefix_sum_offsets_estimated_input_size, decode_retinaclusters, + velo_calculate_phi_and_sort, velo_search_by_triplet, + prefix_sum_offsets_velo_tracks, velo_three_hit_tracks_filter, + prefix_sum_offsets_number_of_three_hit_tracks_filtered, + velo_copy_track_hit_number, + prefix_sum_offsets_velo_track_hit_number, velo_consolidate_tracks) + else: + velo_sequence = Sequence( + host_ut_banks, host_scifi_banks, initialize_lists, velo_banks, + velo_calculate_number_of_candidates, + prefix_sum_offsets_velo_candidates, velo_estimate_input_size, + prefix_sum_offsets_estimated_input_size, velo_masked_clustering, + velo_calculate_phi_and_sort, velo_search_by_triplet, + prefix_sum_offsets_velo_tracks, velo_three_hit_tracks_filter, + prefix_sum_offsets_number_of_three_hit_tracks_filtered, + velo_copy_track_hit_number, + prefix_sum_offsets_velo_track_hit_number, velo_consolidate_tracks) return velo_sequence diff --git a/configuration/sequences/hlt1_pp_retina_cluster.py b/configuration/sequences/hlt1_pp_retina_cluster.py new file mode 100644 index 0000000000000000000000000000000000000000..99d2376d1c6278897f731a5d6094679db9ac64fa --- /dev/null +++ b/configuration/sequences/hlt1_pp_retina_cluster.py @@ -0,0 +1,62 @@ +from definitions.VeloSequence import VeloSequence +from definitions.PVSequence import PVSequence +from definitions.UTSequence import UTSequence +from definitions.ForwardSequence import ForwardSequence +from definitions.MuonSequence import MuonSequence +from definitions.HLT1Sequence import HLT1Sequence +from definitions.algorithms import compose_sequences + +velo_sequence = VeloSequence(retina_decoding=True) + +pv_sequence = PVSequence( + initialize_lists=velo_sequence["initialize_lists"], + velo_copy_track_hit_number=velo_sequence["velo_copy_track_hit_number"], + velo_consolidate_tracks=velo_sequence["velo_consolidate_tracks"], + prefix_sum_offsets_velo_track_hit_number=velo_sequence[ + "prefix_sum_offsets_velo_track_hit_number"]) + +ut_sequence = UTSequence( + initialize_lists=velo_sequence["initialize_lists"], + velo_copy_track_hit_number=velo_sequence["velo_copy_track_hit_number"], + velo_consolidate_tracks=velo_sequence["velo_consolidate_tracks"], + prefix_sum_offsets_velo_track_hit_number=velo_sequence[ + "prefix_sum_offsets_velo_track_hit_number"]) + +forward_sequence = ForwardSequence( + initialize_lists=velo_sequence["initialize_lists"], + velo_copy_track_hit_number=velo_sequence["velo_copy_track_hit_number"], + velo_consolidate_tracks=velo_sequence["velo_consolidate_tracks"], + prefix_sum_offsets_velo_track_hit_number=velo_sequence[ + "prefix_sum_offsets_velo_track_hit_number"], + prefix_sum_ut_tracks=ut_sequence["prefix_sum_ut_tracks"], + prefix_sum_ut_track_hit_number=ut_sequence[ + "prefix_sum_ut_track_hit_number"], + ut_consolidate_tracks=ut_sequence["ut_consolidate_tracks"]) + +muon_sequence = MuonSequence( + initialize_lists=velo_sequence["initialize_lists"], + prefix_sum_forward_tracks=forward_sequence["prefix_sum_forward_tracks"], + prefix_sum_scifi_track_hit_number=forward_sequence[ + "prefix_sum_scifi_track_hit_number"], + scifi_consolidate_tracks_t=forward_sequence["scifi_consolidate_tracks_t"]) + +hlt1_sequence = HLT1Sequence( + initialize_lists=velo_sequence["initialize_lists"], + velo_copy_track_hit_number=velo_sequence["velo_copy_track_hit_number"], + velo_kalman_filter=pv_sequence["velo_kalman_filter"], + prefix_sum_offsets_velo_track_hit_number=velo_sequence[ + "prefix_sum_offsets_velo_track_hit_number"], + pv_beamline_multi_fitter=pv_sequence["pv_beamline_multi_fitter"], + prefix_sum_forward_tracks=forward_sequence["prefix_sum_forward_tracks"], + velo_consolidate_tracks=velo_sequence["velo_consolidate_tracks"], + prefix_sum_ut_tracks=ut_sequence["prefix_sum_ut_tracks"], + prefix_sum_ut_track_hit_number=ut_sequence[ + "prefix_sum_ut_track_hit_number"], + ut_consolidate_tracks=ut_sequence["ut_consolidate_tracks"], + prefix_sum_scifi_track_hit_number=forward_sequence[ + "prefix_sum_scifi_track_hit_number"], + scifi_consolidate_tracks=forward_sequence["scifi_consolidate_tracks_t"], + is_muon=muon_sequence["is_muon_t"]) + +compose_sequences(velo_sequence, pv_sequence, ut_sequence, forward_sequence, + muon_sequence, hlt1_sequence).generate() diff --git a/device/velo/CMakeLists.txt b/device/velo/CMakeLists.txt index 54f4b2c1198c160e93d7ff68795adfbe172c45be..1604e02e4ecc17729c4f17bd901c7610a91993e4 100644 --- a/device/velo/CMakeLists.txt +++ b/device/velo/CMakeLists.txt @@ -1,5 +1,6 @@ file(GLOB velo_prefix_sum "prefix_sum/src/*cu") file(GLOB velo_clustering_sources "mask_clustering/src/*cu") +file(GLOB velo_retinaclusters_sources "retinacluster_decoding/src/*cu") file(GLOB velo_phi_and_sort "calculate_phi_and_sort/src/*cu") file(GLOB velo_search_by_triplet "search_by_triplet/src/*cu") file(GLOB velo_simplified_kalman_filter "simplified_kalman_filter/src/*cu") @@ -9,6 +10,7 @@ include_directories(${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES}) include_directories(common/include) include_directories(prefix_sum/include) include_directories(mask_clustering/include) +include_directories(retinacluster_decoding/include) include_directories(calculate_phi_and_sort/include) include_directories(search_by_triplet/include) include_directories(simplified_kalman_filter/include) @@ -28,6 +30,7 @@ include_directories(${Boost_INCLUDE_DIRS}) allen_add_device_library(Velo STATIC ${velo_clustering_sources} + ${velo_retinaclusters_sources} ${velo_prefix_sum} ${velo_phi_and_sort} ${velo_search_by_triplet} diff --git a/device/velo/common/include/ClusteringDefinitions.cuh b/device/velo/common/include/ClusteringDefinitions.cuh index 44a1a9b4522126c5dbf0f13c53cc254ef252ab92..89f55a71c69d07aaf8f8420b9065dabe6410aeac 100644 --- a/device/velo/common/include/ClusteringDefinitions.cuh +++ b/device/velo/common/include/ClusteringDefinitions.cuh @@ -67,17 +67,17 @@ struct VeloRawEvent { struct VeloRawBank { uint32_t sensor_index; - uint32_t sp_count; - uint32_t* sp_word; - + uint32_t count; + uint32_t* word; + // For MEP format __device__ __host__ VeloRawBank(uint32_t source_id, const char* fragment) { sensor_index = source_id; const char* p = fragment; - sp_count = *((uint32_t*) p); + count = *((uint32_t*) p); p += sizeof(uint32_t); - sp_word = (uint32_t*) p; + word = (uint32_t*) p; } // For Allen format @@ -86,9 +86,9 @@ struct VeloRawBank { const char* p = raw_bank; sensor_index = *((uint32_t*) p); p += sizeof(uint32_t); - sp_count = *((uint32_t*) p); + count = *((uint32_t*) p); p += sizeof(uint32_t); - sp_word = (uint32_t*) p; + word = (uint32_t*) p; } }; diff --git a/device/velo/mask_clustering/src/EstimateInputSize.cu b/device/velo/mask_clustering/src/EstimateInputSize.cu index fb92965f6ea7028d3e638cbac52fa775c2735cd7..016923a3407f0f86c366f3a68f3446993fa390c1 100644 --- a/device/velo/mask_clustering/src/EstimateInputSize.cu +++ b/device/velo/mask_clustering/src/EstimateInputSize.cu @@ -47,8 +47,8 @@ __device__ void estimate_raw_bank_size( { unsigned* estimated_module_pair_size = estimated_input_size + (raw_bank.sensor_index / 8); unsigned found_cluster_candidates = 0; - for (unsigned sp_index = threadIdx.x; sp_index < raw_bank.sp_count; sp_index += blockDim.x) { // Decode sp - const uint32_t sp_word = raw_bank.sp_word[sp_index]; + for (unsigned sp_index = threadIdx.x; sp_index < raw_bank.count; sp_index += blockDim.x) { // Decode sp + const uint32_t sp_word = raw_bank.word[sp_index]; const uint32_t no_sp_neighbours = sp_word & 0x80000000U; const uint32_t sp_addr = (sp_word & 0x007FFF00U) >> 8; const uint8_t sp = sp_word & 0xFFU; @@ -123,8 +123,8 @@ __device__ void estimate_raw_bank_size( const uint32_t sp_row = sp_addr & 0x3FU; const uint32_t sp_col = sp_addr >> 6; - for (unsigned k = 0; k < raw_bank.sp_count; ++k) { - const uint32_t other_sp_word = raw_bank.sp_word[k]; + for (unsigned k = 0; k < raw_bank.count; ++k) { + const uint32_t other_sp_word = raw_bank.word[k]; const uint32_t other_no_sp_neighbours = sp_word & 0x80000000U; if (!other_no_sp_neighbours) { diff --git a/device/velo/mask_clustering/src/MaskedVeloClustering.cu b/device/velo/mask_clustering/src/MaskedVeloClustering.cu index 5a3a28e1eb43c7ce214eb99c9a151a9a42a1d37c..13b04e1fe3757e3fc36b28bc885b3a0293ac1308 100644 --- a/device/velo/mask_clustering/src/MaskedVeloClustering.cu +++ b/device/velo/mask_clustering/src/MaskedVeloClustering.cu @@ -179,9 +179,9 @@ __device__ void no_neighbour_sp( { const float* ltg = g.ltg + g.n_trans * raw_bank.sensor_index; - for (unsigned sp_index = 0; sp_index < raw_bank.sp_count; ++sp_index) { + for (unsigned sp_index = 0; sp_index < raw_bank.count; ++sp_index) { // Decode sp - const uint32_t sp_word = raw_bank.sp_word[sp_index]; + const uint32_t sp_word = raw_bank.word[sp_index]; const uint32_t sp_addr = (sp_word & 0x007FFF00U) >> 8; const uint32_t no_sp_neighbours = sp_word & 0x80000000U; @@ -280,7 +280,7 @@ __device__ void rest_of_clusters( const auto starting_pixel_location = candidate & 0x7; 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_word = raw_bank.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) // This is not a bug @@ -328,8 +328,8 @@ __device__ void rest_of_clusters( // Load SPs // Note: We will pick up the current one, // no need to add a special case - for (unsigned k = 0; k < raw_bank.sp_count; ++k) { - const uint32_t other_sp_word = raw_bank.sp_word[k]; + for (unsigned k = 0; k < raw_bank.count; ++k) { + const uint32_t other_sp_word = raw_bank.word[k]; const uint32_t other_no_sp_neighbours = other_sp_word & 0x80000000U; if (!other_no_sp_neighbours) { const uint32_t other_sp_addr = (other_sp_word & 0x007FFF00U) >> 8; diff --git a/device/velo/mask_clustering/src/VeloCalculateNumberOfCandidates.cu b/device/velo/mask_clustering/src/VeloCalculateNumberOfCandidates.cu index 9356a417bc948d35beaa36283db411c8a055d49a..67b4db1428112cf02e561fb04a7a4a0377870e33 100644 --- a/device/velo/mask_clustering/src/VeloCalculateNumberOfCandidates.cu +++ b/device/velo/mask_clustering/src/VeloCalculateNumberOfCandidates.cu @@ -56,7 +56,7 @@ __global__ void velo_calculate_number_of_candidates::velo_calculate_number_of_ca for (unsigned raw_bank_number = 0; raw_bank_number < raw_event.number_of_raw_banks; ++raw_bank_number) { // Read raw bank const auto raw_bank = VeloRawBank(raw_event.payload + raw_event.raw_bank_offset[raw_bank_number]); - number_of_candidates += raw_bank.sp_count; + number_of_candidates += raw_bank.count; } // The maximum number of candidates is two times the number of SPs @@ -78,7 +78,7 @@ __global__ void velo_calculate_number_of_candidates::velo_calculate_number_of_ca // Create raw bank from MEP layout const auto raw_bank = MEP::raw_bank( parameters.dev_velo_raw_input, parameters.dev_velo_raw_input_offsets, selected_event_number, raw_bank_number); - number_of_candidates += raw_bank.sp_count; + number_of_candidates += raw_bank.count; } parameters.dev_number_of_candidates[event_number] = 2 * number_of_candidates; diff --git a/device/velo/readme.md b/device/velo/readme.md index f8348875b86bf673f70c0c6cd829f55019c974d4..589aa6ae760b7e9db4f4a30459535132b23d9a9e 100644 --- a/device/velo/readme.md +++ b/device/velo/readme.md @@ -18,9 +18,9 @@ raw_bank_offset | uint32_t | number_of_rawbanks ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ sensor_index | uint32_t | 1 | ------------------------------------------------------------------------------ -sp_count | uint32_t | 1 | number_of_rawbanks +count | uint32_t | 1 | number_of_rawbanks ------------------------------------------------------------------------------ -sp_word | uint32_t | sp_count | +word | uint32_t | count | ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ ``` @@ -28,6 +28,5 @@ raw_bank_offset: array containing the offsets to each raw bank, this array is currently filled when running Brunel to create this output; it is needed to process several raw banks in parallel -sp = super pixel -sp_word: contains super pixel address and 8 bit hit pattern +word: contains super pixel address and 8 bit hit pattern diff --git a/device/velo/retinacluster_decoding/include/CalculateNumberOfRetinaClustersPerSensor.cuh b/device/velo/retinacluster_decoding/include/CalculateNumberOfRetinaClustersPerSensor.cuh new file mode 100644 index 0000000000000000000000000000000000000000..2f60f65ab5d8a8d7d750147bc1b836660260e072 --- /dev/null +++ b/device/velo/retinacluster_decoding/include/CalculateNumberOfRetinaClustersPerSensor.cuh @@ -0,0 +1,38 @@ +#pragma once + +#include "DeviceAlgorithm.cuh" +#include "ClusteringDefinitions.cuh" + +namespace calculate_number_of_retinaclusters_each_sensor { + DEFINE_PARAMETERS( + Parameters, + (HOST_INPUT(host_number_of_selected_events_t, uint), host_number_of_selected_events), + (DEVICE_INPUT(dev_event_list_t, uint), dev_event_list), + (DEVICE_INPUT(dev_velo_retina_raw_input_t, char), dev_velo_retina_raw_input), + (DEVICE_INPUT(dev_velo_retina_raw_input_offsets_t, uint), dev_velo_retina_raw_input_offsets), + (DEVICE_OUTPUT(dev_each_sensor_size_t, uint), dev_each_sensor_size), + (PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions), block_dim_prop)) + + __global__ void calculate_number_of_retinaclusters_each_sensor(Parameters parameters); + + __global__ void calculate_number_of_retinaclusters_each_sensor_mep(Parameters parameters); + + struct calculate_number_of_retinaclusters_each_sensor_t : public DeviceAlgorithm, Parameters { + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const; + + void operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants&, + HostBuffers&, + cudaStream_t& cuda_stream, + cudaEvent_t&) const; + + private: + Property m_block_dim {this, {{256, 1, 1}}}; + }; +} // namespace calculate_number_of_retinaclusters_each_sensor diff --git a/device/velo/retinacluster_decoding/include/DecodeRetinaClusters.cuh b/device/velo/retinacluster_decoding/include/DecodeRetinaClusters.cuh new file mode 100644 index 0000000000000000000000000000000000000000..e88322ab1532bf5113dc55028128272d24a361fc --- /dev/null +++ b/device/velo/retinacluster_decoding/include/DecodeRetinaClusters.cuh @@ -0,0 +1,49 @@ +#pragma once + +#include +#include +#include "ClusteringDefinitions.cuh" +#include "VeloEventModel.cuh" +#include "DeviceAlgorithm.cuh" + +namespace decode_retinaclusters { + DEFINE_PARAMETERS( + Parameters, + (HOST_INPUT(host_total_number_of_velo_clusters_t, uint), host_total_number_of_velo_clusters), + (HOST_INPUT(host_number_of_selected_events_t, uint), host_number_of_selected_events), + (DEVICE_INPUT(dev_velo_retina_raw_input_t, char), dev_velo_retina_raw_input), + (DEVICE_INPUT(dev_velo_retina_raw_input_offsets_t, uint), dev_velo_retina_raw_input_offsets), + (DEVICE_INPUT(dev_offsets_each_sensor_size_t, uint), dev_offsets_each_sensor_size), + (DEVICE_INPUT(dev_event_list_t, uint), dev_event_list), + (DEVICE_OUTPUT(dev_module_cluster_num_t, uint), dev_module_pair_cluster_num), + (DEVICE_OUTPUT(dev_offsets_module_pair_cluster_t, uint), dev_offsets_module_pair_cluster), + (DEVICE_OUTPUT(dev_velo_cluster_container_t, char), dev_velo_cluster_container), + (PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions), block_dim_prop)) + + __global__ void decode_retinaclusters( + Parameters parameters, + const VeloGeometry* dev_velo_geometry); + + __global__ void decode_retinaclusters_mep( + Parameters parameters, + const VeloGeometry* dev_velo_geometry); + + struct decode_retinaclusters_t : public DeviceAlgorithm, Parameters { + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const; + + void operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants&, + HostBuffers&, + cudaStream_t& cuda_stream, + cudaEvent_t&) const; + + private: + Property m_block_dim {this, {{8, 32, 1}}}; + }; +} // namespace decode_retinaclusters diff --git a/device/velo/retinacluster_decoding/src/CalculateNumberOfRetinaClustersPerSensor.cu b/device/velo/retinacluster_decoding/src/CalculateNumberOfRetinaClustersPerSensor.cu new file mode 100644 index 0000000000000000000000000000000000000000..1507fdb3418edbda0e42444c0dec1a1c27061b89 --- /dev/null +++ b/device/velo/retinacluster_decoding/src/CalculateNumberOfRetinaClustersPerSensor.cu @@ -0,0 +1,73 @@ +#include +#include + +void calculate_number_of_retinaclusters_each_sensor::calculate_number_of_retinaclusters_each_sensor_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const +{ + if (logger::verbosity() >= logger::debug) { + debug_cout << "# of events = " << first(arguments) << std::endl; + } + + set_size( + arguments, first(arguments) * Velo::Constants::n_modules * Velo::Constants::n_sensors_per_module); +} + +void calculate_number_of_retinaclusters_each_sensor::calculate_number_of_retinaclusters_each_sensor_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants&, + HostBuffers&, + cudaStream_t& cuda_stream, + cudaEvent_t&) const +{ + initialize(arguments, 0, cuda_stream); + + if (runtime_options.mep_layout) { + global_function(calculate_number_of_retinaclusters_each_sensor_mep)(dim3(first(arguments)), property(), cuda_stream)(arguments); + } + else { + global_function(calculate_number_of_retinaclusters_each_sensor)( + dim3(first(arguments)), property(), cuda_stream)(arguments); + } +} + +__global__ void calculate_number_of_retinaclusters_each_sensor::calculate_number_of_retinaclusters_each_sensor(calculate_number_of_retinaclusters_each_sensor::Parameters parameters) +{ + const auto event_number = blockIdx.x; + const auto selected_event_number = parameters.dev_event_list[event_number]; + + const char* raw_input = parameters.dev_velo_retina_raw_input + parameters.dev_velo_retina_raw_input_offsets[selected_event_number]; + uint* each_sensor_size = parameters.dev_each_sensor_size + event_number * Velo::Constants::n_modules * Velo::Constants::n_sensors_per_module; + + // Read raw event + const auto raw_event = VeloRawEvent(raw_input); + + for (uint raw_bank_number = threadIdx.x; raw_bank_number < raw_event.number_of_raw_banks; + raw_bank_number += blockDim.x) { + // Read raw bank + const auto raw_bank = VeloRawBank(raw_event.payload + raw_event.raw_bank_offset[raw_bank_number]); + each_sensor_size[raw_bank.sensor_index] = raw_bank.count; + } + +} +__global__ void calculate_number_of_retinaclusters_each_sensor::calculate_number_of_retinaclusters_each_sensor_mep(calculate_number_of_retinaclusters_each_sensor::Parameters parameters) +{ + const uint event_number = blockIdx.x; + const uint selected_event_number = parameters.dev_event_list[event_number]; + + uint* each_sensor_size = parameters.dev_each_sensor_size + event_number * Velo::Constants::n_modules * Velo::Constants::n_sensors_per_module; + + // Read raw event + auto const number_of_raw_banks = parameters.dev_velo_retina_raw_input_offsets[0]; + + for (uint raw_bank_number = threadIdx.x; raw_bank_number < number_of_raw_banks; raw_bank_number += blockDim.x) { + + // Create raw bank from MEP layout + const auto raw_bank = MEP::raw_bank( + parameters.dev_velo_retina_raw_input, parameters.dev_velo_retina_raw_input_offsets, selected_event_number, raw_bank_number); + each_sensor_size[raw_bank.sensor_index] = raw_bank.count; + } +} diff --git a/device/velo/retinacluster_decoding/src/DecodeRetinaClusters.cu b/device/velo/retinacluster_decoding/src/DecodeRetinaClusters.cu new file mode 100644 index 0000000000000000000000000000000000000000..9658a9197d6d8053f3e9f59b42e2b58714d7062e --- /dev/null +++ b/device/velo/retinacluster_decoding/src/DecodeRetinaClusters.cu @@ -0,0 +1,170 @@ +#include +#include + +void decode_retinaclusters::decode_retinaclusters_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const +{ + set_size( + arguments, first(arguments) * Velo::Constants::n_module_pairs); + set_size(arguments, + first(arguments) * Velo::Clusters::element_size); + set_size(arguments, + first(arguments) * Velo::Constants::n_module_pairs + 1); +} + +void decode_retinaclusters::decode_retinaclusters_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + HostBuffers&, + cudaStream_t& cuda_stream, + cudaEvent_t&) const +{ + initialize(arguments, 0, cuda_stream); + initialize(arguments, 0, cuda_stream); + + if (runtime_options.mep_layout) { + global_function(decode_retinaclusters_mep)( + dim3(first(arguments)), property(), cuda_stream)(arguments, constants.dev_velo_geometry); + } else { + global_function(decode_retinaclusters)( + dim3(first(arguments)), property(), cuda_stream)(arguments, constants.dev_velo_geometry); + } +} + +__device__ void put_retinaclusters_into_container( + Velo::Clusters velo_cluster_container, + VeloGeometry const& g, + uint const cluster_start, + VeloRawBank const& raw_bank) +{ + const float* ltg = g.ltg + g.n_trans * raw_bank.sensor_index; + + for (uint rc_index = threadIdx.y; rc_index < raw_bank.count; rc_index += blockDim.y) { + // Decode cluster + const uint32_t word = raw_bank.word[rc_index]; + + const uint32_t cx = ( word >> 14 ) & 0x3FF; + const float fx = ( ( word >> 11 ) & 0x7 ) / 8.f; + const uint32_t cy = ( word >> 3 ) & 0xFF; + const float fy = ( word & 0x7FF ) / 8.f; + + const uint32_t chip = cx >> VP::ChipColumns_division; + const float local_x = g.local_x[cx] + fx * g.x_pitch[cx]; + const float local_y = (0.5f + fy) * Velo::Constants::pixel_size; + + const float gx = ( ltg[0] * local_x + ltg[1] * local_y + ltg[9] ); + const float gy = ( ltg[3] * local_x + ltg[4] * local_y + ltg[10] ); + const float gz = ( ltg[6] * local_x + ltg[7] * local_y + ltg[11] ); + const uint cid = get_channel_id(raw_bank.sensor_index, chip, cx & VP::ChipColumns_mask, cy); + + velo_cluster_container.set_x(cluster_start + rc_index, gx); + velo_cluster_container.set_y(cluster_start + rc_index, gy); + velo_cluster_container.set_z(cluster_start + rc_index, gz); + velo_cluster_container.set_id(cluster_start + rc_index, get_lhcb_id(cid)); + + } +} + + +__global__ void decode_retinaclusters::decode_retinaclusters( + decode_retinaclusters::Parameters parameters, + const VeloGeometry* dev_velo_geometry) +{ + const uint number_of_events = gridDim.x; + const uint event_number = blockIdx.x; + const uint selected_event_number = parameters.dev_event_list[event_number]; + + const char* raw_input = parameters.dev_velo_retina_raw_input + parameters.dev_velo_retina_raw_input_offsets[selected_event_number]; + const uint* sensor_cluster_start = + parameters.dev_offsets_each_sensor_size + event_number * Velo::Constants::n_modules * Velo::Constants::n_sensors_per_module; + uint* module_pair_cluster_num = parameters.dev_module_pair_cluster_num + event_number * Velo::Constants::n_module_pairs; + uint* offsets_pair_module_size = parameters.dev_offsets_module_pair_cluster + 1 + event_number * Velo::Constants::n_module_pairs; + uint* first_number_of_dev_offsets_module_pair_cluster = parameters.dev_offsets_module_pair_cluster; + first_number_of_dev_offsets_module_pair_cluster[0] = 0; + + const uint* offsets_each_sensor_size = parameters.dev_offsets_each_sensor_size + event_number * Velo::Constants::n_modules * Velo::Constants::n_sensors_per_module; + + // Local pointers to parameters.dev_velo_cluster_container + const uint estimated_number_of_clusters = + parameters.dev_offsets_each_sensor_size[Velo::Constants::n_module_pairs * number_of_events * 8]; + auto velo_cluster_container = Velo::Clusters {parameters.dev_velo_cluster_container, estimated_number_of_clusters}; + + // Load Velo geometry (assume it is the same for all events) + const VeloGeometry& g = *dev_velo_geometry; + + + // Read raw event + const auto raw_event = VeloRawEvent(raw_input); + + for (uint raw_bank_number = threadIdx.x; raw_bank_number < raw_event.number_of_raw_banks; + raw_bank_number += blockDim.x) { + const auto module_pair_number = raw_bank_number / 8; + const auto offsets_sensor_behind = (module_pair_number + 1) * 8; + const auto offsets_sensor_in_front = module_pair_number * 8; + const uint cluster_start = sensor_cluster_start[raw_bank_number]; + offsets_pair_module_size[module_pair_number] = offsets_each_sensor_size[offsets_sensor_behind]; + module_pair_cluster_num[module_pair_number] = offsets_each_sensor_size[offsets_sensor_behind] - offsets_each_sensor_size[offsets_sensor_in_front]; + + // Read raw bank + const auto raw_bank = VeloRawBank(raw_event.payload + raw_event.raw_bank_offset[raw_bank_number]); + put_retinaclusters_into_container( + velo_cluster_container, + g, + cluster_start, + raw_bank); + } + +} + +__global__ void decode_retinaclusters::decode_retinaclusters_mep( + decode_retinaclusters::Parameters parameters, + const VeloGeometry* dev_velo_geometry) +{ + const uint number_of_events = gridDim.x; + const uint event_number = blockIdx.x; + const uint selected_event_number = parameters.dev_event_list[event_number]; + + const uint* sensor_cluster_start = + parameters.dev_offsets_each_sensor_size + event_number * Velo::Constants::n_modules * Velo::Constants::n_sensors_per_module; + uint* module_pair_cluster_num = parameters.dev_module_pair_cluster_num + event_number * Velo::Constants::n_module_pairs; + uint* offsets_pair_module_size = parameters.dev_offsets_module_pair_cluster + 1 + event_number * Velo::Constants::n_module_pairs; + uint* first_number_of_dev_offsets_module_pair_cluster = parameters.dev_offsets_module_pair_cluster; + first_number_of_dev_offsets_module_pair_cluster[0] = 0; + + const uint* offsets_each_sensor_size = parameters.dev_offsets_each_sensor_size + event_number * Velo::Constants::n_modules * Velo::Constants::n_sensors_per_module; + + // Local pointers to parameters.dev_velo_cluster_container + const uint estimated_number_of_clusters = + parameters.dev_offsets_each_sensor_size[Velo::Constants::n_module_pairs * number_of_events * 8]; + auto velo_cluster_container = Velo::Clusters {parameters.dev_velo_cluster_container, estimated_number_of_clusters}; + + + // Load Velo geometry (assume it is the same for all events) + const VeloGeometry& g = *dev_velo_geometry; + + // Read raw event + auto const number_of_raw_banks = parameters.dev_velo_retina_raw_input_offsets[0]; + + for (uint raw_bank_number = threadIdx.x; raw_bank_number < number_of_raw_banks; raw_bank_number += blockDim.x) { + const auto module_pair_number = raw_bank_number / 8; + const auto offsets_sensor_behind = (module_pair_number + 1) * 8; + const auto offsets_sensor_in_front = module_pair_number * 8; + const uint cluster_start = sensor_cluster_start[raw_bank_number]; + offsets_pair_module_size[module_pair_number] = offsets_each_sensor_size[offsets_sensor_behind]; + module_pair_cluster_num[module_pair_number] = offsets_each_sensor_size[offsets_sensor_behind] - offsets_each_sensor_size[offsets_sensor_in_front]; + + // Read raw bank + const auto raw_bank = MEP::raw_bank( + parameters.dev_velo_retina_raw_input, parameters.dev_velo_retina_raw_input_offsets, selected_event_number, raw_bank_number); + put_retinaclusters_into_container( + velo_cluster_container, + g, + cluster_start, + raw_bank); + } + +} diff --git a/host/velo/clustering/src/Clustering.cpp b/host/velo/clustering/src/Clustering.cpp index 980fb9e45c1691ccf5223b3b6dd36a0bc91ca33d..a449c1e2c9ddc078f65e6bf633d4fba555e7d3e8 100644 --- a/host/velo/clustering/src/Clustering.cpp +++ b/host/velo/clustering/src/Clustering.cpp @@ -33,8 +33,8 @@ std::vector> clustering( // const unsigned int module = sensor / Velo::Constants::n_sensors_per_module; // const float* ltg = g.ltg + 16 * sensor; - for (unsigned int j = 0; j < velo_raw_bank.sp_count; ++j) { - const uint32_t sp_word = *(velo_raw_bank.sp_word + j); + for (unsigned int j = 0; j < velo_raw_bank.count; ++j) { + const uint32_t sp_word = *(velo_raw_bank.word + j); uint8_t sp = sp_word & 0xFFU; diff --git a/main/include/BankTypes.h b/main/include/BankTypes.h index eadd164f308149b45461a90aaf2e578e25c175c3..6581a2bdb330f81e712dc39f1ed0ec1528dae1f0 100644 --- a/main/include/BankTypes.h +++ b/main/include/BankTypes.h @@ -13,13 +13,14 @@ namespace { using gsl::span; } -constexpr auto NBankTypes = 10; -enum class BankTypes { VP, UT, FT, MUON, ODIN, Rich, ECal, HCal, OTRaw, OTError, Unknown }; +constexpr auto NBankTypes = 11; +enum class BankTypes { VP, VPRetinaCluster, UT, FT, MUON, ODIN, Rich, ECal, HCal, OTRaw, OTError, Unknown }; // Average size of all raw banks of a given type per // subdetector, in kB, measured in simulated minbias events. // FIXME: make this configurable const std::unordered_map BankSizes = {{BankTypes::VP, 12.f}, + {BankTypes::VPRetinaCluster, 12.f}, {BankTypes::UT, 7.f}, {BankTypes::FT, 9.f}, {BankTypes::MUON, 1.2f}, diff --git a/main/src/Allen.cpp b/main/src/Allen.cpp index 006839302d96351d2df43345a35fe5681e05500f..db0df4d5da971e2b08535924e901dd2d9f71ff93 100644 --- a/main/src/Allen.cpp +++ b/main/src/Allen.cpp @@ -336,6 +336,7 @@ extern "C" int allen( // Raw data input folders const auto folder_name_velopix_raw = folder_data + folder_rawdata + "VP"; + const auto folder_name_veloretina_raw = folder_data + folder_rawdata + "VPRetinaCluster"; const auto folder_name_UT_raw = folder_data + folder_rawdata + "UT"; const auto folder_name_SciFi_raw = folder_data + folder_rawdata + "FTCluster"; const auto folder_name_Muon_raw = folder_data + folder_rawdata + "Muon"; @@ -379,9 +380,13 @@ extern "C" int allen( non_stop, // Run the application non-stop !mep_layout, // MEPs should be transposed to Allen layout receivers}; // Map of receiver to MPI rank to receive from - input_provider = - std::make_unique>( - number_of_slices, *events_per_slice, n_events, split_string(mep_input, ","), config); + input_provider = std::make_unique>(number_of_slices, *events_per_slice, n_events, split_string(mep_input, ","), config); } else if (!mdf_input.empty()) { mep_layout = false; @@ -391,18 +396,30 @@ extern "C" int allen( *events_per_slice * 10 + 1, // mximum number event of offsets in read buffer *events_per_slice, // number of events per read buffer n_io_reps}; // number of loops over the input files - input_provider = - std::make_unique>( - number_of_slices, *events_per_slice, n_events, split_string(mdf_input, ","), config); + input_provider = std::make_unique>(number_of_slices, *events_per_slice, n_events, split_string(mdf_input, ","), config); } else { mep_layout = false; // The binary input provider expects the folders for the bank types as connections - std::vector connections = { - folder_name_velopix_raw, folder_name_UT_raw, folder_name_SciFi_raw, folder_name_Muon_raw, folder_name_ODIN_raw}; - input_provider = - std::make_unique>( - number_of_slices, *events_per_slice, n_events, std::move(connections), n_io_reps, file_list); + std::vector connections = {folder_name_velopix_raw, + folder_name_veloretina_raw, + folder_name_UT_raw, + folder_name_SciFi_raw, + folder_name_Muon_raw, + folder_name_ODIN_raw}; + input_provider = std::make_unique>(number_of_slices, *events_per_slice, n_events, std::move(connections), n_io_reps, file_list); } // Load constant parameters from JSON diff --git a/main/src/BankTypes.cpp b/main/src/BankTypes.cpp index 843cc90259383b596d36a6614888ecd0bfa7b873..8f62ac8f1e1115ae7e59ce39031c3dacfcd13c79 100644 --- a/main/src/BankTypes.cpp +++ b/main/src/BankTypes.cpp @@ -6,6 +6,7 @@ namespace { const std::map BankNames = {{BankTypes::VP, "VP"}, + {BankTypes::VPRetinaCluster, "VPRetinaCluster"}, {BankTypes::UT, "UT"}, {BankTypes::FT, "FTCluster"}, {BankTypes::MUON, "Muon"}, diff --git a/main/src/CheckVP.cpp b/main/src/CheckVP.cpp index c13386059e9b07f1567dc6af751ffa4572a37e07..d7e68ebe898953906c3adbfc3156b42f8c448506 100644 --- a/main/src/CheckVP.cpp +++ b/main/src/CheckVP.cpp @@ -24,13 +24,13 @@ bool check_velopix_events(const std::vector& events, const std::vector bank_types = {{LHCb::RawBank::VP, BankTypes::VP}, - {LHCb::RawBank::UT, BankTypes::UT}, - {LHCb::RawBank::FTCluster, BankTypes::FT}, - {LHCb::RawBank::Muon, BankTypes::MUON}, - {LHCb::RawBank::ODIN, BankTypes::ODIN}}; + const std::unordered_map bank_types = { + {LHCb::RawBank::VP, BankTypes::VP}, + {LHCb::RawBank::VPRetinaCluster, BankTypes::VPRetinaCluster}, + {LHCb::RawBank::UT, BankTypes::UT}, + {LHCb::RawBank::FTCluster, BankTypes::FT}, + {LHCb::RawBank::Muon, BankTypes::MUON}, + {LHCb::RawBank::ODIN, BankTypes::ODIN}}; using buffer_map = std::unordered_map, std::vector>>; diff --git a/mdf/test/test_providers.cpp b/mdf/test/test_providers.cpp index 9389a7b83cc8f036e8129522afd37bed6edc6a0a..dc835a189f8dfc49bbdb5b27b9c876ac292e86a6 100644 --- a/mdf/test/test_providers.cpp +++ b/mdf/test/test_providers.cpp @@ -34,8 +34,22 @@ namespace { Config s_config; MDFProviderConfig mdf_config {true, 2, 1}; - unique_ptr> mdf; - unique_ptr> binary; + unique_ptr> + mdf; + unique_ptr> + binary; size_t slice_mdf = 0, slice_binary = 0; size_t filled_mdf = 0, filled_binary = 0; @@ -118,20 +132,31 @@ int main(int argc, char* argv[]) } } } - for (auto sd : {"UT"s, "VP"s, "FTCluster"s, "Muon"s, "ODIN"s}) { + for (auto sd : {"UT"s, "VP"s, "VPRetinaCluster"s, "FTCluster"s, "Muon"s, "ODIN"s}) { s_config.banks_dirs.push_back(directory + "/banks/" + sd); } if (s_config.run) { // Allocate providers and get slices - mdf = make_unique>( - s_config.n_slices, s_config.n_events, s_config.n_events, s_config.mdf_files, mdf_config); + mdf = make_unique>(s_config.n_slices, s_config.n_events, s_config.n_events, s_config.mdf_files, mdf_config); bool good = false, timed_out = false, done = false; std::tie(good, done, timed_out, slice_mdf, filled_mdf) = mdf->get_slice(); auto const& events_mdf = mdf->event_ids(slice_mdf); - binary = make_unique>( + binary = make_unique>( s_config.n_slices, s_config.n_events, s_config.n_events, @@ -154,8 +179,13 @@ int main(int argc, char* argv[]) auto size_fun = [pf](BankTypes) -> std::tuple { return {std::lround(average_event_size * pf * bank_size_fudge_factor * kB), pf}; }; - mep_slices = - allocate_slices(1, size_fun); + mep_slices = allocate_slices< + BankTypes::VP, + BankTypes::VPRetinaCluster, + BankTypes::UT, + BankTypes::FT, + BankTypes::MUON, + BankTypes::ODIN>(1, size_fun); transpose_mep(mep_slices, 0, mep_header, mep_span, s_config.n_events); } @@ -205,6 +235,7 @@ TEMPLATE_TEST_CASE( "MDF versus Binary", "[MDF binary]", BTTag, + BTTag, BTTag, BTTag, BTTag) @@ -245,6 +276,7 @@ TEMPLATE_TEST_CASE( "Binary vs MEP", "[MEP binary]", BTTag, + BTTag, BTTag, BTTag, BTTag) diff --git a/test/reference/bsphiphi_geforcertx2080ti.txt b/test/reference/bsphiphi_geforcertx2080ti_hlt1_pp_default.txt similarity index 100% rename from test/reference/bsphiphi_geforcertx2080ti.txt rename to test/reference/bsphiphi_geforcertx2080ti_hlt1_pp_default.txt diff --git a/test/reference/bsphiphi_geforcertx2080ti_hlt1_pp_retina_cluster.txt b/test/reference/bsphiphi_geforcertx2080ti_hlt1_pp_retina_cluster.txt new file mode 100644 index 0000000000000000000000000000000000000000..474a78dcfe8990427b60cb637a58d0054fa2e891 --- /dev/null +++ b/test/reference/bsphiphi_geforcertx2080ti_hlt1_pp_retina_cluster.txt @@ -0,0 +1,94 @@ +Velo tracks: +TrackChecker output : 4574/ 285569 1.60% ghosts +01_velo : 119427/ 121547 98.26% ( 98.30%), 2844 ( 2.33%) clones, pur 99.42%, hit eff 96.01% +02_long : 68304/ 68861 99.19% ( 99.24%), 1162 ( 1.67%) clones, pur 99.51%, hit eff 97.18% +03_long_P>5GeV : 43134/ 43397 99.39% ( 99.44%), 612 ( 1.40%) clones, pur 99.53%, hit eff 97.69% +04_long_strange : 3044/ 3127 97.35% ( 97.66%), 57 ( 1.84%) clones, pur 99.06%, hit eff 96.55% +05_long_strange_P>5GeV : 1429/ 1468 97.34% ( 97.52%), 22 ( 1.52%) clones, pur 98.85%, hit eff 97.40% +06_long_fromB : 4395/ 4453 98.70% ( 98.84%), 61 ( 1.37%) clones, pur 99.25%, hit eff 97.05% +07_long_fromB_P>5GeV : 3639/ 3678 98.94% ( 98.88%), 44 ( 1.19%) clones, pur 99.24%, hit eff 97.36% +08_long_electrons : 4818/ 5253 91.72% ( 91.89%), 145 ( 2.92%) clones, pur 95.56%, hit eff 93.34% +09_long_fromB_electrons : 217/ 233 93.13% ( 94.50%), 14 ( 6.06%) clones, pur 96.28%, hit eff 93.21% +10_long_fromB_electrons_P>5GeV : 153/ 162 94.44% ( 96.62%), 9 ( 5.56%) clones, pur 96.84%, hit eff 94.82% + + +Primary vertices: +REC and MC vertices matched by dz distance +MC PV is reconstructible if at least 4 tracks are reconstructed +MC PV is isolated if dz to closest reconstructible MC PV > 10.00 mm +REC and MC vertices matched by dz distance + +All : 0.925 ( 5327/ 5758) +Isolated : 0.965 ( 2881/ 2984) +Close : 0.882 ( 2446/ 2774) +False rate : 0.014 ( 73/ 5400) +Real false rate : 0.014 ( 73/ 5400) +Clones : 0.000 ( 0/ 5327) + + +Velo+UT tracks: +TrackChecker output : 4036/ 58914 6.85% ghosts +01_velo : 52206/ 121547 42.95% ( 43.21%), 533 ( 1.01%) clones, pur 99.30%, hit eff 96.68% +02_velo+UT : 52087/ 105844 49.21% ( 49.53%), 532 ( 1.01%) clones, pur 99.32%, hit eff 96.68% +03_velo+UT_P>5GeV : 37286/ 52476 71.05% ( 71.47%), 360 ( 0.96%) clones, pur 99.41%, hit eff 97.53% +04_velo+notLong : 10657/ 52686 20.23% ( 20.25%), 102 ( 0.95%) clones, pur 99.11%, hit eff 95.47% +05_velo+UT+notLong : 10556/ 37859 27.88% ( 27.99%), 102 ( 0.96%) clones, pur 99.20%, hit eff 95.45% +06_velo+UT+notLong_P>5GeV : 5923/ 9814 60.35% ( 60.88%), 43 ( 0.72%) clones, pur 99.41%, hit eff 97.43% +07_long : 41549/ 68861 60.34% ( 60.77%), 431 ( 1.03%) clones, pur 99.35%, hit eff 96.99% +08_long_P>5GeV : 31381/ 43397 72.31% ( 72.78%), 318 ( 1.00%) clones, pur 99.40%, hit eff 97.55% +09_long_fromB : 3697/ 4453 83.02% ( 85.66%), 35 ( 0.94%) clones, pur 99.26%, hit eff 97.49% +10_long_fromB_P>5GeV : 3314/ 3678 90.10% ( 91.38%), 29 ( 0.87%) clones, pur 99.27%, hit eff 97.62% +11_long_electrons : 895/ 5253 17.04% ( 17.89%), 23 ( 2.51%) clones, pur 96.61%, hit eff 95.07% +12_long_fromB_electrons : 95/ 233 40.77% ( 45.26%), 3 ( 3.06%) clones, pur 97.87%, hit eff 95.49% +13_long_fromB_electrons_P>5GeV : 90/ 162 55.56% ( 60.40%), 3 ( 3.23%) clones, pur 97.92%, hit eff 95.25% + + +Forward tracks: +TrackChecker output : 1374/ 38721 3.55% ghosts +for P>3GeV,Pt>0.5GeV : 752/ 24989 3.01% ghosts +01_long : 35270/ 68861 51.22% ( 52.23%), 360 ( 1.01%) clones, pur 99.05%, hit eff 95.65% +02_long_P>5GeV : 29337/ 43397 67.60% ( 68.43%), 294 ( 0.99%) clones, pur 99.18%, hit eff 96.61% +03_long_strange : 1073/ 3127 34.31% ( 34.52%), 12 ( 1.11%) clones, pur 98.66%, hit eff 94.99% +04_long_strange_P>5GeV : 788/ 1468 53.68% ( 54.24%), 7 ( 0.88%) clones, pur 98.78%, hit eff 96.43% +05_long_fromB : 3380/ 4453 75.90% ( 78.49%), 31 ( 0.91%) clones, pur 99.18%, hit eff 96.82% +06_long_fromB_P>5GeV : 3152/ 3678 85.70% ( 86.91%), 28 ( 0.88%) clones, pur 99.22%, hit eff 97.10% +07_long_electrons : 679/ 5253 12.93% ( 13.48%), 21 ( 3.00%) clones, pur 97.83%, hit eff 95.35% +08_long_electrons_P>5GeV : 588/ 2619 22.45% ( 23.87%), 18 ( 2.97%) clones, pur 97.91%, hit eff 96.22% +09_long_fromB_electrons : 84/ 233 36.05% ( 39.90%), 3 ( 3.45%) clones, pur 98.34%, hit eff 95.71% +10_long_fromB_electrons_P>5GeV : 79/ 162 48.77% ( 52.51%), 3 ( 3.66%) clones, pur 98.50%, hit eff 96.15% + + +Muon matching: +Muon fraction in all MCPs: 13985/ 1035270 0.01% +Muon fraction in MCPs to which a track(s) was matched: 456/ 43111 0.01% +Correctly identified muons with isMuon: 373/ 456 81.80% +Correctly identified muons from strange decays with isMuon: 0/ 0 -nan% +Correctly identified muons from B decays with isMuon: 92/ 109 84.40% +Tracks identified as muon with isMuon, but matched to non-muon MCP: 4118/ 42655 9.65% +Ghost tracks identified as muon with isMuon: 170/ 1374 12.37% + + +HLT1 rates: +ErrorEvent: 0/ 1000, ( 0.00 +/- 0.00) kHz +PassThrough: 1000/ 1000, (30000.00 +/- 0.00) kHz +NoBeams: 1/ 1000, ( 30.00 +/- 29.98) kHz +BeamOne: 0/ 1000, ( 0.00 +/- 0.00) kHz +BeamTwo: 0/ 1000, ( 0.00 +/- 0.00) kHz +BothBeams: 0/ 1000, ( 0.00 +/- 0.00) kHz +ODINNoBias: 0/ 1000, ( 0.00 +/- 0.00) kHz +ODINLumi: 0/ 1000, ( 0.00 +/- 0.00) kHz +GECPassthrough: 1000/ 1000, (30000.00 +/- 0.00) kHz +VeloMicroBias: 1/ 1000, ( 30.00 +/- 29.98) kHz +TrackMVA: 244/ 1000, ( 7320.00 +/- 407.45) kHz +TrackMuonMVA: 8/ 1000, ( 240.00 +/- 84.51) kHz +SingleHighPtMuon: 3/ 1000, ( 90.00 +/- 51.88) kHz +LowPtMuon: 130/ 1000, ( 3900.00 +/- 319.05) kHz +TwoTrackMVA: 427/ 1000, (12810.00 +/- 469.26) kHz +DiMuonHighMass: 13/ 1000, ( 390.00 +/- 107.46) kHz +DiMuonLowMass: 30/ 1000, ( 900.00 +/- 161.83) kHz +LowPtDiMuon: 46/ 1000, ( 1380.00 +/- 198.73) kHz +DiMuonSoft: 0/ 1000, ( 0.00 +/- 0.00) kHz +D2KPi: 38/ 1000, ( 1140.00 +/- 181.38) kHz +D2PiPi: 36/ 1000, ( 1080.00 +/- 176.73) kHz +D2KK: 16/ 1000, ( 480.00 +/- 119.04) kHz +Inclusive: 1000/ 1000, (30000.00 +/- 0.00) kHz diff --git a/test/reference/bsphiphi_x862630v4.txt b/test/reference/bsphiphi_x862630v4_hlt1_pp_default.txt similarity index 100% rename from test/reference/bsphiphi_x862630v4.txt rename to test/reference/bsphiphi_x862630v4_hlt1_pp_default.txt