diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index b9f0e080910..94bac79414f 100644 --- a/.gitlab-ci.yml +++ b/.gitlab-ci.yml @@ -186,6 +186,43 @@ test_exatrkx_python: - pytest -rFsv -k torch --collect-only - pytest -rFsv -k gpu-torch # For now only test torch GPU pipeline +build_gnn_tensorrt: + stage: build + image: ghcr.io/acts-project/ubuntu2404_tensorrt:74 + variables: + DEPENDENCY_URL: https://acts.web.cern.ch/ACTS/ci/ubuntu-24.04/deps.$DEPENDENCY_TAG.tar.zst + + cache: + key: ccache-${CI_JOB_NAME}-${CI_COMMIT_REF_SLUG}-${CCACHE_KEY_SUFFIX} + fallback_keys: + - ccache-${CI_JOB_NAME}-${CI_DEFAULT_BRANCH}-${CCACHE_KEY_SUFFIX} + when: always + paths: + - ${CCACHE_DIR} + + tags: + - docker-gpu-nvidia + + script: + - git clone $CLONE_URL src + - cd src + - git checkout $HEAD_SHA + - source CI/dependencies.sh + - cd .. + - mkdir build + - > + cmake -B build -S src + -DACTS_BUILD_PLUGIN_EXATRKX=ON + -DACTS_EXATRKX_ENABLE_TORCH=OFF + -DACTS_EXATRKX_ENABLE_CUDA=ON + -DACTS_EXATRKX_ENABLE_TENSORRT=ON + -DPython_EXECUTABLE=$(which python3) + -DCMAKE_CUDA_ARCHITECTURES="75;86" + - ccache -z + - cmake --build build -- -j6 + - ccache -s + + build_linux_ubuntu: stage: build image: registry.cern.ch/ghcr.io/acts-project/ubuntu2404:63 diff --git a/CMakeLists.txt b/CMakeLists.txt index 5487b731292..a0dad0cf109 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -435,13 +435,6 @@ if(ACTS_BUILD_PLUGIN_EXATRKX) else() message(STATUS "Build Exa.TrkX plugin for CPU only") endif() - if(NOT (ACTS_EXATRKX_ENABLE_ONNX OR ACTS_EXATRKX_ENABLE_TORCH)) - message( - FATAL_ERROR - "When building the Exa.TrkX plugin, at least one of ACTS_EXATRKX_ENABLE_ONNX \ - and ACTS_EXATRKX_ENABLE_TORCHSCRIPT must be enabled." - ) - endif() if(ACTS_EXATRKX_ENABLE_TORCH) find_package(TorchScatter REQUIRED) endif() diff --git a/Examples/Python/src/ExaTrkXTrackFinding.cpp b/Examples/Python/src/ExaTrkXTrackFinding.cpp index 5ef30a17df3..1eaac52c027 100644 --- a/Examples/Python/src/ExaTrkXTrackFinding.cpp +++ b/Examples/Python/src/ExaTrkXTrackFinding.cpp @@ -11,6 +11,7 @@ #include "Acts/Plugins/ExaTrkX/ExaTrkXPipeline.hpp" #include "Acts/Plugins/ExaTrkX/OnnxEdgeClassifier.hpp" #include "Acts/Plugins/ExaTrkX/OnnxMetricLearning.hpp" +#include "Acts/Plugins/ExaTrkX/TensorRTEdgeClassifier.hpp" #include "Acts/Plugins/ExaTrkX/TorchEdgeClassifier.hpp" #include "Acts/Plugins/ExaTrkX/TorchMetricLearning.hpp" #include "Acts/Plugins/ExaTrkX/TorchTruthGraphMetricsHook.hpp" @@ -113,6 +114,31 @@ void addExaTrkXTrackFinding(Context &ctx) { } #endif +#ifdef ACTS_EXATRKX_WITH_TENSORRT + { + using Alg = Acts::TensorRTEdgeClassifier; + using Config = Alg::Config; + + auto alg = + py::class_>( + mex, "TensorRTEdgeClassifier") + .def(py::init([](const Config &c, Logging::Level lvl) { + return std::make_shared( + c, getDefaultLogger("EdgeClassifier", lvl)); + }), + py::arg("config"), py::arg("level")) + .def_property_readonly("config", &Alg::config); + + auto c = py::class_(alg, "Config").def(py::init<>()); + ACTS_PYTHON_STRUCT_BEGIN(c, Config); + ACTS_PYTHON_MEMBER(modelPath); + ACTS_PYTHON_MEMBER(selectedFeatures); + ACTS_PYTHON_MEMBER(cut); + ACTS_PYTHON_MEMBER(deviceID); + ACTS_PYTHON_MEMBER(doSigmoid); + } +#endif + #ifdef ACTS_EXATRKX_WITH_CUDA { using Alg = Acts::CudaTrackBuilding; diff --git a/Plugins/ExaTrkX/CMakeLists.txt b/Plugins/ExaTrkX/CMakeLists.txt index e70a896b8d5..b192f4938fc 100644 --- a/Plugins/ExaTrkX/CMakeLists.txt +++ b/Plugins/ExaTrkX/CMakeLists.txt @@ -24,6 +24,20 @@ if(ACTS_EXATRKX_ENABLE_TORCH) ) endif() +if(ACTS_EXATRKX_ENABLE_TENSORRT) + find_package(TensorRT REQUIRED) + message(STATUS "Found TensorRT ${TensorRT_VERSION}") + target_link_libraries( + ActsPluginExaTrkX + PUBLIC trt::nvinfer trt::nvinfer_plugin + ) + target_sources(ActsPluginExaTrkX PRIVATE src/TensorRTEdgeClassifier.cpp) + target_compile_definitions( + ActsPluginExaTrkX + PUBLIC ACTS_EXATRKX_WITH_TENSORRT + ) +endif() + target_include_directories( ActsPluginExaTrkX PUBLIC diff --git a/Plugins/ExaTrkX/include/Acts/Plugins/ExaTrkX/TensorRTEdgeClassifier.hpp b/Plugins/ExaTrkX/include/Acts/Plugins/ExaTrkX/TensorRTEdgeClassifier.hpp new file mode 100644 index 00000000000..dd3ac90625f --- /dev/null +++ b/Plugins/ExaTrkX/include/Acts/Plugins/ExaTrkX/TensorRTEdgeClassifier.hpp @@ -0,0 +1,63 @@ +// This file is part of the ACTS project. +// +// Copyright (C) 2016 CERN for the benefit of the ACTS project +// +// This Source Code Form is subject to the terms of the Mozilla Public +// License, v. 2.0. If a copy of the MPL was not distributed with this +// file, You can obtain one at https://mozilla.org/MPL/2.0/. + +#pragma once + +#include "Acts/Plugins/ExaTrkX/Stages.hpp" +#include "Acts/Utilities/Logger.hpp" + +#include +#include + +#include + +namespace nvinfer1 { +class IRuntime; +class ICudaEngine; +class ILogger; +class IExecutionContext; +} // namespace nvinfer1 + +namespace Acts { + +class TensorRTEdgeClassifier final : public Acts::EdgeClassificationBase { + public: + struct Config { + std::string modelPath; + std::vector selectedFeatures = {}; + float cut = 0.5; + + std::size_t numExecutionContexts = 1; + }; + + TensorRTEdgeClassifier(const Config &cfg, + std::unique_ptr logger); + ~TensorRTEdgeClassifier(); + + std::tuple operator()( + std::any nodeFeatures, std::any edgeIndex, std::any edgeFeatures = {}, + const ExecutionContext &execContext = {}) override; + + Config config() const { return m_cfg; } + torch::Device device() const override { return torch::kCUDA; }; + + private: + std::unique_ptr m_logger; + const auto &logger() const { return *m_logger; } + + Config m_cfg; + + std::unique_ptr m_runtime; + std::unique_ptr m_engine; + std::unique_ptr m_trtLogger; + + mutable std::mutex m_contextMutex; + mutable std::vector> m_contexts; +}; + +} // namespace Acts diff --git a/Plugins/ExaTrkX/src/TensorRTEdgeClassifier.cpp b/Plugins/ExaTrkX/src/TensorRTEdgeClassifier.cpp new file mode 100644 index 00000000000..5fd0903afcf --- /dev/null +++ b/Plugins/ExaTrkX/src/TensorRTEdgeClassifier.cpp @@ -0,0 +1,208 @@ +// This file is part of the ACTS project. +// +// Copyright (C) 2016 CERN for the benefit of the ACTS project +// +// This Source Code Form is subject to the terms of the Mozilla Public +// License, v. 2.0. If a copy of the MPL was not distributed with this +// file, You can obtain one at https://mozilla.org/MPL/2.0/. + +#include "Acts/Plugins/ExaTrkX/TensorRTEdgeClassifier.hpp" + +#include "Acts/Plugins/ExaTrkX/detail/CudaUtils.cuh" +#include "Acts/Plugins/ExaTrkX/detail/Utils.hpp" + +#include +#include +#include + +#include +#include +#include +#include +#include + +#include "printCudaMemInfo.hpp" + +using namespace torch::indexing; + +namespace { + +class TensorRTLogger : public nvinfer1::ILogger { + std::unique_ptr m_logger; + + public: + TensorRTLogger(Acts::Logging::Level lvl) + : m_logger(Acts::getDefaultLogger("TensorRT", lvl)) {} + + void log(Severity severity, const char *msg) noexcept override { + const auto &logger = *m_logger; + switch (severity) { + case Severity::kVERBOSE: + ACTS_DEBUG(msg); + break; + case Severity::kINFO: + ACTS_INFO(msg); + break; + case Severity::kWARNING: + ACTS_WARNING(msg); + break; + case Severity::kERROR: + ACTS_ERROR(msg); + break; + case Severity::kINTERNAL_ERROR: + ACTS_FATAL(msg); + break; + } + } +}; + +} // namespace + +namespace Acts { + +TensorRTEdgeClassifier::TensorRTEdgeClassifier( + const Config &cfg, std::unique_ptr _logger) + : m_logger(std::move(_logger)), + m_cfg(cfg), + m_trtLogger(std::make_unique(m_logger->level())) { + auto status = initLibNvInferPlugins(m_trtLogger.get(), ""); + if (!status) { + throw std::runtime_error("Failed to initialize TensorRT plugins"); + } + + std::size_t fsize = + std::filesystem::file_size(std::filesystem::path(m_cfg.modelPath)); + std::vector engineData(fsize); + + ACTS_DEBUG("Load '" << m_cfg.modelPath << "' with size " << fsize); + + std::ifstream engineFile(m_cfg.modelPath); + if (!engineFile) { + throw std::runtime_error("Failed to open engine file"); + } else if (!engineFile.read(engineData.data(), fsize)) { + throw std::runtime_error("Failed to read engine file"); + } + + m_runtime.reset(nvinfer1::createInferRuntime(*m_trtLogger)); + if (!m_runtime) { + throw std::runtime_error("Failed to create TensorRT runtime"); + } + + m_engine.reset(m_runtime->deserializeCudaEngine(engineData.data(), fsize)); + if (!m_engine) { + throw std::runtime_error("Failed to deserialize CUDA engine"); + } + + for (auto i = 0ul; i < m_cfg.numExecutionContexts; ++i) { + ACTS_DEBUG("Create execution context " << i); + m_contexts.emplace_back(m_engine->createExecutionContext()); + if (!m_contexts.back()) { + throw std::runtime_error("Failed to create execution context"); + } + } + + std::size_t freeMem, totalMem; + cudaMemGetInfo(&freeMem, &totalMem); + ACTS_DEBUG("Used CUDA memory after TensorRT initialization: " + << (totalMem - freeMem) * 1e-9 << " / " << totalMem * 1e-9 + << " GB"); +} + +TensorRTEdgeClassifier::~TensorRTEdgeClassifier() {} + +std::tuple +TensorRTEdgeClassifier::operator()(std::any inNodeFeatures, + std::any inEdgeIndex, + std::any inEdgeFeatures, + const ExecutionContext &execContext) { + assert(execContext.device.is_cuda()); + decltype(std::chrono::high_resolution_clock::now()) t0, t1, t2, t3, t4; + t0 = std::chrono::high_resolution_clock::now(); + + c10::cuda::CUDAStreamGuard(execContext.stream.value()); + + auto nodeFeatures = + std::any_cast(inNodeFeatures).to(execContext.device); + + auto edgeIndex = + std::any_cast(inEdgeIndex).to(execContext.device); + ACTS_DEBUG("edgeIndex: " << detail::TensorDetails{edgeIndex}); + + auto edgeFeatures = + std::any_cast(inEdgeFeatures).to(execContext.device); + ACTS_DEBUG("edgeFeatures: " << detail::TensorDetails{edgeFeatures}); + + t1 = std::chrono::high_resolution_clock::now(); + + // get a context from the list of contexts + std::unique_ptr context; + while (context == nullptr) { + std::lock_guard lock(m_contextMutex); + if (!m_contexts.empty()) { + context = std::move(m_contexts.back()); + m_contexts.pop_back(); + } + } + assert(context != nullptr); + + context->setInputShape( + "x", nvinfer1::Dims2{nodeFeatures.size(0), nodeFeatures.size(1)}); + context->setTensorAddress("x", nodeFeatures.data_ptr()); + + context->setInputShape("edge_index", + nvinfer1::Dims2{edgeIndex.size(0), edgeIndex.size(1)}); + context->setTensorAddress("edge_index", edgeIndex.data_ptr()); + + context->setInputShape( + "edge_attr", nvinfer1::Dims2{edgeFeatures.size(0), edgeFeatures.size(1)}); + context->setTensorAddress("edge_attr", edgeFeatures.data_ptr()); + + auto scores = torch::empty( + edgeIndex.size(1), + torch::TensorOptions().device(torch::kCUDA).dtype(torch::kFloat32)); + context->setTensorAddress("output", scores.data_ptr()); + + t2 = std::chrono::high_resolution_clock::now(); + + auto stream = execContext.stream.value().stream(); + auto status = context->enqueueV3(stream); + if (!status) { + throw std::runtime_error("Failed to execute TensorRT model"); + } + ACTS_CUDA_CHECK(cudaStreamSynchronize(stream)); + + t3 = std::chrono::high_resolution_clock::now(); + + { + std::lock_guard lock(m_contextMutex); + m_contexts.push_back(std::move(context)); + } + + scores.sigmoid_(); + + ACTS_VERBOSE("Size after classifier: " << scores.size(0)); + ACTS_VERBOSE("Slice of classified output:\n" + << scores.slice(/*dim=*/0, /*start=*/0, /*end=*/9)); + printCudaMemInfo(logger()); + + torch::Tensor mask = scores > m_cfg.cut; + torch::Tensor edgesAfterCut = edgeIndex.index({Slice(), mask}); + + scores = scores.masked_select(mask); + ACTS_VERBOSE("Size after score cut: " << edgesAfterCut.size(1)); + printCudaMemInfo(logger()); + + t4 = std::chrono::high_resolution_clock::now(); + + auto milliseconds = [](const auto &a, const auto &b) { + return std::chrono::duration(b - a).count(); + }; + ACTS_DEBUG("Time anycast: " << milliseconds(t0, t1)); + ACTS_DEBUG("Time alloc, set shape " << milliseconds(t1, t2)); + ACTS_DEBUG("Time inference: " << milliseconds(t2, t3)); + ACTS_DEBUG("Time sigmoid and cut: " << milliseconds(t3, t4)); + + return {nodeFeatures, edgesAfterCut, edgeFeatures, scores}; +} + +} // namespace Acts diff --git a/cmake/FindTensorRT.cmake b/cmake/FindTensorRT.cmake new file mode 100644 index 00000000000..42994b7a3f2 --- /dev/null +++ b/cmake/FindTensorRT.cmake @@ -0,0 +1,182 @@ +# ~~~ +# Copyright 2021 Olivier Le Doeuff +# Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: +# The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. +# This module defines the following variables: +# +# - TensorRT_FOUND: A boolean specifying whether or not TensorRT was found. +# - TensorRT_VERSION: The exact version of TensorRT found +# - TensorRT_VERSION_MAJOR: The major version of TensorRT. +# - TensorRT_VERSION_MINOR: The minor version of TensorRT. +# - TensorRT_VERSION_PATCH: The patch version of TensorRT. +# - TensorRT_VERSION_TWEAK: The tweak version of TensorRT. +# - TensorRT_INCLUDE_DIRS: The path to TensorRT ``include`` folder containing the header files required to compile a project linking against TensorRT. +# - TensorRT_LIBRARY_DIRS: The path to TensorRT library directory that contains libraries. +# +# This module create following targets: +# - trt::nvinfer +# - trt::nvinfer_plugin +# - trt::nvonnxparser +# - trt::nvparsers +# This script was inspired from https://github.com/NicolasIRAGNE/CMakeScripts +# This script was inspired from https://github.com/NVIDIA/tensorrt-laboratory/blob/master/cmake/FindTensorRT.cmake +# +# Hints +# ^^^^^ +# A user may set ``TensorRT_ROOT`` to an installation root to tell this module where to look. +# ~~~ + +if(NOT TensorRT_FIND_COMPONENTS) + set(TensorRT_FIND_COMPONENTS nvinfer nvinfer_plugin nvonnxparser nvparsers) +endif() +set(TensorRT_LIBRARIES) + +# find the include directory of TensorRT +find_path( + TensorRT_INCLUDE_DIR + NAMES NvInfer.h + PATHS ${TensorRT_ROOT} + ENV TensorRT_ROOT + PATH_SUFFIXES include +) + +string(FIND ${TensorRT_INCLUDE_DIR} "NOTFOUND" _include_dir_notfound) +if(NOT _include_dir_notfound EQUAL -1) + if(TensorRT_FIND_REQUIRED) + message( + FATAL_ERROR + "Fail to find TensorRT, please set TensorRT_ROOT. Include path not found." + ) + endif() + return() +endif() +set(TensorRT_INCLUDE_DIRS ${TensorRT_INCLUDE_DIR}) + +# Extract version of tensorrt +if(EXISTS "${TensorRT_INCLUDE_DIR}/NvInferVersion.h") + file( + STRINGS + "${TensorRT_INCLUDE_DIR}/NvInferVersion.h" + TensorRT_MAJOR + REGEX "^#define NV_TENSORRT_MAJOR [0-9]+.*$" + ) + file( + STRINGS + "${TensorRT_INCLUDE_DIR}/NvInferVersion.h" + TensorRT_MINOR + REGEX "^#define NV_TENSORRT_MINOR [0-9]+.*$" + ) + file( + STRINGS + "${TensorRT_INCLUDE_DIR}/NvInferVersion.h" + TensorRT_PATCH + REGEX "^#define NV_TENSORRT_PATCH [0-9]+.*$" + ) + file( + STRINGS + "${TensorRT_INCLUDE_DIR}/NvInferVersion.h" + TensorRT_TWEAK + REGEX "^#define NV_TENSORRT_BUILD [0-9]+.*$" + ) + + string( + REGEX REPLACE + "^#define NV_TENSORRT_MAJOR ([0-9]+).*$" + "\\1" + TensorRT_VERSION_MAJOR + "${TensorRT_MAJOR}" + ) + string( + REGEX REPLACE + "^#define NV_TENSORRT_MINOR ([0-9]+).*$" + "\\1" + TensorRT_VERSION_MINOR + "${TensorRT_MINOR}" + ) + string( + REGEX REPLACE + "^#define NV_TENSORRT_PATCH ([0-9]+).*$" + "\\1" + TensorRT_VERSION_PATCH + "${TensorRT_PATCH}" + ) + string( + REGEX REPLACE + "^#define NV_TENSORRT_BUILD ([0-9]+).*$" + "\\1" + TensorRT_VERSION_TWEAK + "${TensorRT_TWEAK}" + ) + set(TensorRT_VERSION + "${TensorRT_VERSION_MAJOR}.${TensorRT_VERSION_MINOR}.${TensorRT_VERSION_PATCH}.${TensorRT_VERSION_TWEAK}" + ) +endif() + +function(_find_trt_component component) + # Find library for component (ie nvinfer, nvparsers, etc...) + find_library( + TensorRT_${component}_LIBRARY + NAMES ${component} + PATHS ${TensorRT_ROOT} ${TENSORRT_LIBRARY_DIR} + ENV TensorRT_ROOT + ) + + string(FIND ${TensorRT_${component}_LIBRARY} "NOTFOUND" _library_not_found) + + if(NOT TensorRT_LIBRARY_DIR) + get_filename_component(_path ${TensorRT_${component}_LIBRARY} DIRECTORY) + set(TensorRT_LIBRARY_DIR + "${_path}" + CACHE INTERNAL + "TensorRT_LIBRARY_DIR" + ) + endif() + + if(NOT TensorRT_LIBRARY_DIRS) + get_filename_component(_path ${TensorRT_${component}_LIBRARY} DIRECTORY) + set(TensorRT_LIBRARY_DIRS + "${_path}" + CACHE INTERNAL + "TensorRT_LIBRARY_DIRS" + ) + endif() + + # Library found, and doesn't already exists + if(_library_not_found EQUAL -1 AND NOT TARGET trt::${component}) + set(TensorRT_${component}_FOUND + TRUE + CACHE INTERNAL + "Found ${component}" + ) + + # Create a target + add_library(trt::${component} IMPORTED INTERFACE) + target_include_directories( + trt::${component} + SYSTEM + INTERFACE "${TensorRT_INCLUDE_DIRS}" + ) + target_link_libraries( + trt::${component} + INTERFACE "${TensorRT_${component}_LIBRARY}" + ) + set(TensorRT_LIBRARIES + ${TensorRT_LIBRARIES} + ${TensorRT_${component}_LIBRARY} + ) + endif() +endfunction() + +# Find each components +foreach(component IN LISTS TensorRT_FIND_COMPONENTS) + _find_trt_component(${component}) +endforeach() + +include(FindPackageHandleStandardArgs) +find_package_handle_standard_args( + TensorRT + HANDLE_COMPONENTS + VERSION_VAR TensorRT_VERSION + REQUIRED_VARS TensorRT_INCLUDE_DIR +)