From 3f93292782077c61bb93a2c4816ac0e12d7d5e84 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 11 Dec 2019 00:00:50 +0100 Subject: [PATCH 1/7] Add back dictionary for vector --- DataFormats/Math/src/classes_def.xml | 1 + 1 file changed, 1 insertion(+) diff --git a/DataFormats/Math/src/classes_def.xml b/DataFormats/Math/src/classes_def.xml index 12aeda101983e..51b3abbab4790 100755 --- a/DataFormats/Math/src/classes_def.xml +++ b/DataFormats/Math/src/classes_def.xml @@ -10,6 +10,7 @@ + From 86cec95a63511e90ca5015212d809741ec5888a0 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 10 Dec 2019 17:55:30 +0100 Subject: [PATCH 2/7] Exercises for the 2019 Patatrack tutorial --- Patatrack/Tutorial/plugins/BuildFile.xml | 6 ++ .../plugins/CompareCartesianVectors.cc | 84 +++++++++++++++++++ .../plugins/ConvertToCartesianVectors.cc | 65 ++++++++++++++ .../plugins/ConvertToCartesianVectorsCUDA.cc | 70 ++++++++++++++++ .../plugins/GenerateCylindricalVectors.cc | 63 ++++++++++++++ .../Tutorial/plugins/PrintCartesianVectors.cc | 47 +++++++++++ .../plugins/PrintCylindricalVectors.cc | 48 +++++++++++ Patatrack/Tutorial/plugins/cudavectors.cu | 27 ++++++ Patatrack/Tutorial/plugins/cudavectors.h | 22 +++++ .../test/benchmarkCartesianVectors.py | 23 +++++ .../test/benchmarkCartesianVectorsCUDA.py | 27 ++++++ .../Tutorial/test/compareCartesianVectors.py | 37 ++++++++ .../test/generateCylindricalVectors.py | 30 +++++++ .../Tutorial/test/printCartesianVectors.py | 27 ++++++ .../test/printCartesianVectorsCUDA.py | 31 +++++++ .../Tutorial/test/printCylindricalVectors.py | 23 +++++ 16 files changed, 630 insertions(+) create mode 100644 Patatrack/Tutorial/plugins/BuildFile.xml create mode 100644 Patatrack/Tutorial/plugins/CompareCartesianVectors.cc create mode 100644 Patatrack/Tutorial/plugins/ConvertToCartesianVectors.cc create mode 100644 Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc create mode 100644 Patatrack/Tutorial/plugins/GenerateCylindricalVectors.cc create mode 100644 Patatrack/Tutorial/plugins/PrintCartesianVectors.cc create mode 100644 Patatrack/Tutorial/plugins/PrintCylindricalVectors.cc create mode 100644 Patatrack/Tutorial/plugins/cudavectors.cu create mode 100644 Patatrack/Tutorial/plugins/cudavectors.h create mode 100644 Patatrack/Tutorial/test/benchmarkCartesianVectors.py create mode 100644 Patatrack/Tutorial/test/benchmarkCartesianVectorsCUDA.py create mode 100644 Patatrack/Tutorial/test/compareCartesianVectors.py create mode 100644 Patatrack/Tutorial/test/generateCylindricalVectors.py create mode 100644 Patatrack/Tutorial/test/printCartesianVectors.py create mode 100644 Patatrack/Tutorial/test/printCartesianVectorsCUDA.py create mode 100644 Patatrack/Tutorial/test/printCylindricalVectors.py diff --git a/Patatrack/Tutorial/plugins/BuildFile.xml b/Patatrack/Tutorial/plugins/BuildFile.xml new file mode 100644 index 0000000000000..5e307e5c27188 --- /dev/null +++ b/Patatrack/Tutorial/plugins/BuildFile.xml @@ -0,0 +1,6 @@ + + + + + + diff --git a/Patatrack/Tutorial/plugins/CompareCartesianVectors.cc b/Patatrack/Tutorial/plugins/CompareCartesianVectors.cc new file mode 100644 index 0000000000000..6d1ef6b6a5d5c --- /dev/null +++ b/Patatrack/Tutorial/plugins/CompareCartesianVectors.cc @@ -0,0 +1,84 @@ +// system include files +#include +#include +#include +#include +#include + +// CMSSW include files +#include "DataFormats/Math/interface/Vector3D.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/Frameworkfwd.h" +#include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/Framework/interface/one/EDAnalyzer.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/InputTag.h" + +class CompareCartesianVectors : public edm::one::EDAnalyzer<> { +public: + explicit CompareCartesianVectors(const edm::ParameterSet&); + ~CompareCartesianVectors() = default; + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + +private: + using CartesianVectors = std::vector; + + virtual void analyze(const edm::Event&, const edm::EventSetup&) override; + + bool compare(math::XYZVectorF const& first, math::XYZVectorF const& second) const; + + edm::EDGetTokenT first_; + edm::EDGetTokenT second_; + const double precision_; +}; + +CompareCartesianVectors::CompareCartesianVectors(const edm::ParameterSet& config) + : first_(consumes(config.getParameter("first"))), + second_(consumes(config.getParameter("second"))), + precision_(config.getParameter("precision")) {} + +void CompareCartesianVectors::analyze(const edm::Event& event, const edm::EventSetup& setup) { + auto const& first = event.get(first_); + auto const& second = event.get(second_); + if (first.size() != second.size()) { + std::cout << "The two collections have different size: " << first.size() << " and " << second.size() << "." + << std::endl; + return; + } + bool consistent = true; + for (unsigned int i = 0; i < first.size(); ++i) { + if (not compare(first[i], second[i])) { + if (consistent) { + std::cout << "Found inconsistent elements:" << std::endl; + consistent = false; + } + std::cout << std::setprecision(9) << "(" << first[i].x() << ", " << first[i].y() << ", " << first[i].z() + << ") vs (" << second[i].x() << ", " << second[i].y() << ", " << second[i].z() << ")" << std::endl; + } + } + if (consistent) { + std::cout << "All elements are consistent within " << precision_ << std::endl; + } +} + +bool CompareCartesianVectors::compare(math::XYZVectorF const& first, math::XYZVectorF const& second) const { + if (std::abs(first.x() - second.x()) > std::abs(first.x() + second.x()) * precision_) + return false; + if (std::abs(first.y() - second.y()) > std::abs(first.y() + second.y()) * precision_) + return false; + if (std::abs(first.z() - second.z()) > std::abs(first.z() + second.z()) * precision_) + return false; + return true; +} + +void CompareCartesianVectors::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + desc.add("first", edm::InputTag("firstCartesianVectors")); + desc.add("second", edm::InputTag("secondCartesianVectors")); + desc.add("precision", 1.e-6); + descriptions.addWithDefaultLabel(desc); +} + +// define this as a plug-in +DEFINE_FWK_MODULE(CompareCartesianVectors); diff --git a/Patatrack/Tutorial/plugins/ConvertToCartesianVectors.cc b/Patatrack/Tutorial/plugins/ConvertToCartesianVectors.cc new file mode 100644 index 0000000000000..b7abd0a81a6d5 --- /dev/null +++ b/Patatrack/Tutorial/plugins/ConvertToCartesianVectors.cc @@ -0,0 +1,65 @@ +// system include files +#include +#include +#include + +// CMSSW include files +#include "DataFormats/Math/interface/Vector3D.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/Frameworkfwd.h" +#include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/Framework/interface/stream/EDProducer.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/StreamID.h" + +class ConvertToCartesianVectors : public edm::stream::EDProducer<> { +public: + explicit ConvertToCartesianVectors(const edm::ParameterSet&); + ~ConvertToCartesianVectors() = default; + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + +private: + using CartesianVectors = std::vector; + using CylindricalVectors = std::vector; + + static void convert(math::RhoEtaPhiVectorF const& cilindrical, math::XYZVectorF & cartesian); + + virtual void produce(edm::Event&, const edm::EventSetup&) override; + + edm::EDGetTokenT input_; + edm::EDPutTokenT output_; +}; + +ConvertToCartesianVectors::ConvertToCartesianVectors(const edm::ParameterSet& config) + : input_(consumes(config.getParameter("input"))) { + output_ = produces(); +} + +void ConvertToCartesianVectors::convert(math::RhoEtaPhiVectorF const& cilindrical, math::XYZVectorF & cartesian) { + cartesian.SetCoordinates(cilindrical.rho() * std::cos(cilindrical.phi()), + cilindrical.rho() * std::sin(cilindrical.phi()), + cilindrical.rho() * std::sinh(cilindrical.eta())); +} + +void ConvertToCartesianVectors::produce(edm::Event& event, const edm::EventSetup& setup) { + auto const& input = event.get(input_); + auto elements = input.size(); + auto product = std::make_unique(elements); + + // convert the vectors from cylindrical to cartesian coordinates + for (unsigned int i = 0; i < elements; ++i) { + convert(input[i], (*product)[i]); + } + + event.put(output_, std::move(product)); +} + +void ConvertToCartesianVectors::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + desc.add("input", edm::InputTag("cylindricalVectors")); + descriptions.addWithDefaultLabel(desc); +} + +// define this as a plug-in +DEFINE_FWK_MODULE(ConvertToCartesianVectors); diff --git a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc new file mode 100644 index 0000000000000..598a10a69a32e --- /dev/null +++ b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc @@ -0,0 +1,70 @@ +// system include files +#include +#include +#include + +// CMSSW include files +#include "DataFormats/Math/interface/Vector3D.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/Frameworkfwd.h" +#include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/Framework/interface/stream/EDProducer.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/StreamID.h" +#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" + +#include "cudavectors.h" + +class ConvertToCartesianVectorsCUDA : public edm::stream::EDProducer<> { +public: + explicit ConvertToCartesianVectorsCUDA(const edm::ParameterSet&); + ~ConvertToCartesianVectorsCUDA() = default; + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + +private: + using CartesianVectors = std::vector; + using CylindricalVectors = std::vector; + + virtual void produce(edm::Event&, const edm::EventSetup&) override; + + edm::EDGetTokenT input_; + edm::EDPutTokenT output_; +}; + +ConvertToCartesianVectorsCUDA::ConvertToCartesianVectorsCUDA(const edm::ParameterSet& config) + : input_(consumes(config.getParameter("input"))) { + output_ = produces(); +} + +void ConvertToCartesianVectorsCUDA::produce(edm::Event& event, const edm::EventSetup& setup) { + auto const& input = event.get(input_); + auto elements = input.size(); + auto product = std::make_unique(elements); + + // allocate memory on the GPU for the cylindrical and cartesian vectors + // fill here ... + + // copy the input data to the GPU + // fill here ... + + // convert the vectors from cylindrical to cartesian coordinates, on the GPU + // fill here ... + + // copy the result from the GPU + // fill here ... + + // free the GPU memory + // fill here ... + + event.put(output_, std::move(product)); +} + +void ConvertToCartesianVectorsCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + desc.add("input", edm::InputTag("cylindricalVectors")); + descriptions.addWithDefaultLabel(desc); +} + +// define this as a plug-in +DEFINE_FWK_MODULE(ConvertToCartesianVectorsCUDA); diff --git a/Patatrack/Tutorial/plugins/GenerateCylindricalVectors.cc b/Patatrack/Tutorial/plugins/GenerateCylindricalVectors.cc new file mode 100644 index 0000000000000..2b798f543c767 --- /dev/null +++ b/Patatrack/Tutorial/plugins/GenerateCylindricalVectors.cc @@ -0,0 +1,63 @@ +// system include files +#include +#include +#include + +// CMSSW include files +#include "DataFormats/Math/interface/Vector3D.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/Frameworkfwd.h" +#include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/Framework/interface/stream/EDProducer.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/StreamID.h" + +class GenerateCylindricalVectors : public edm::stream::EDProducer<> { +public: + explicit GenerateCylindricalVectors(const edm::ParameterSet&); + ~GenerateCylindricalVectors() = default; + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + +private: + using CylindricalVectors = std::vector; + + void beginStream(edm::StreamID) override {} + void produce(edm::Event&, const edm::EventSetup&) override; + void endStream() override {} + + std::mt19937 generator_; + std::lognormal_distribution genPt_; + std::uniform_real_distribution genEta_; + std::uniform_real_distribution genPhi_; + const uint32_t size_; + + edm::EDPutTokenT output_; +}; + +GenerateCylindricalVectors::GenerateCylindricalVectors(const edm::ParameterSet& config) + : generator_(std::random_device()()), + genPt_(3, 0.6), + genEta_(-5., +5.), + genPhi_(0., 2 * M_PI), + size_(config.getParameter("size")) // number of CylindricalVectors to generate +{ + output_ = produces(); +} + +void GenerateCylindricalVectors::produce(edm::Event& event, const edm::EventSetup& setup) { + auto product = std::make_unique(size_); + for (auto& p : *product) { + p.SetCoordinates(genPt_(generator_), genEta_(generator_), genPhi_(generator_)); + } + event.put(output_, std::move(product)); +} + +void GenerateCylindricalVectors::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + desc.add("size", 1000)->setComment("number of generated elements"); + descriptions.addWithDefaultLabel(desc); +} + +// define this as a plug-in +DEFINE_FWK_MODULE(GenerateCylindricalVectors); diff --git a/Patatrack/Tutorial/plugins/PrintCartesianVectors.cc b/Patatrack/Tutorial/plugins/PrintCartesianVectors.cc new file mode 100644 index 0000000000000..237f4dc31dcf4 --- /dev/null +++ b/Patatrack/Tutorial/plugins/PrintCartesianVectors.cc @@ -0,0 +1,47 @@ +// system include files +#include +#include + +// CMSSW include files +#include "DataFormats/Math/interface/Vector3D.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/Frameworkfwd.h" +#include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/Framework/interface/one/EDAnalyzer.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/InputTag.h" + +class PrintCartesianVectors : public edm::one::EDAnalyzer<> { +public: + explicit PrintCartesianVectors(const edm::ParameterSet&); + ~PrintCartesianVectors() = default; + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + +private: + using CartesianVectors = std::vector; + + virtual void analyze(const edm::Event&, const edm::EventSetup&) override; + + edm::EDGetTokenT input_; +}; + +PrintCartesianVectors::PrintCartesianVectors(const edm::ParameterSet& config) + : input_(consumes(config.getParameter("input"))) {} + +void PrintCartesianVectors::analyze(const edm::Event& event, const edm::EventSetup& setup) { + for (auto const& v : event.get(input_)) { + std::cout << std::fixed << "x: " << std::setw(6) << std::setprecision(2) << v.x() << ", y: " << std::setw(6) + << std::setprecision(2) << v.y() << ", z: " << std::setw(8) << std::setprecision(2) << v.z() << std::endl; + } + std::cout << std::endl; +} + +void PrintCartesianVectors::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + desc.add("input", edm::InputTag("cartesianVectors")); + descriptions.addWithDefaultLabel(desc); +} + +// define this as a plug-in +DEFINE_FWK_MODULE(PrintCartesianVectors); diff --git a/Patatrack/Tutorial/plugins/PrintCylindricalVectors.cc b/Patatrack/Tutorial/plugins/PrintCylindricalVectors.cc new file mode 100644 index 0000000000000..a5da4705a769c --- /dev/null +++ b/Patatrack/Tutorial/plugins/PrintCylindricalVectors.cc @@ -0,0 +1,48 @@ +// system include files +#include +#include + +// CMSSW include files +#include "DataFormats/Math/interface/Vector3D.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/Frameworkfwd.h" +#include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/Framework/interface/one/EDAnalyzer.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/InputTag.h" + +class PrintCylindricalVectors : public edm::one::EDAnalyzer<> { +public: + explicit PrintCylindricalVectors(const edm::ParameterSet&); + ~PrintCylindricalVectors() = default; + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + +private: + using CylindricalVectors = std::vector; + + virtual void analyze(const edm::Event&, const edm::EventSetup&) override; + + edm::EDGetTokenT input_; +}; + +PrintCylindricalVectors::PrintCylindricalVectors(const edm::ParameterSet& config) + : input_(consumes(config.getParameter("input"))) {} + +void PrintCylindricalVectors::analyze(const edm::Event& event, const edm::EventSetup& setup) { + for (auto const& v : event.get(input_)) { + std::cout << std::fixed << "pT: " << std::setw(6) << std::setprecision(2) << v.rho() << ", eta: " << std::setw(6) + << std::setprecision(2) << v.eta() << ", phi: " << std::setw(6) << std::setprecision(2) << v.phi() + << std::endl; + } + std::cout << std::endl; +} + +void PrintCylindricalVectors::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + desc.add("input", edm::InputTag("cylindricalVectors")); + descriptions.addWithDefaultLabel(desc); +} + +// define this as a plug-in +DEFINE_FWK_MODULE(PrintCylindricalVectors); diff --git a/Patatrack/Tutorial/plugins/cudavectors.cu b/Patatrack/Tutorial/plugins/cudavectors.cu new file mode 100644 index 0000000000000..8da23582aad2c --- /dev/null +++ b/Patatrack/Tutorial/plugins/cudavectors.cu @@ -0,0 +1,27 @@ +// system include files +#include + +// CUDA include files +#include + +// CMSSW include files +#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" +#include "cudavectors.h" + +namespace cudavectors { + + __host__ __device__ inline void convert(CylindricalVector const& cylindrical, CartesianVector & cartesian) { + // fill here ... + } + + __global__ void convertKernel(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size) { + // fill here ... + } + + void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size) { + // fill here ... + //convertKernel<<>>(cylindrical, cartesian, size); + cudaCheck(cudaGetLastError()); + } + +} // namespace cudavectors diff --git a/Patatrack/Tutorial/plugins/cudavectors.h b/Patatrack/Tutorial/plugins/cudavectors.h new file mode 100644 index 0000000000000..c20625f37b2b0 --- /dev/null +++ b/Patatrack/Tutorial/plugins/cudavectors.h @@ -0,0 +1,22 @@ +#ifndef cudavectors_h +#define cudavectors_h + +namespace cudavectors { + + struct CylindricalVector { + float rho; + float eta; + float phi; + }; + + struct CartesianVector { + float x; + float y; + float z; + }; + + void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size); + +} // namespace cudavectors + +#endif // cudavectors_h diff --git a/Patatrack/Tutorial/test/benchmarkCartesianVectors.py b/Patatrack/Tutorial/test/benchmarkCartesianVectors.py new file mode 100644 index 0000000000000..fd7e829140603 --- /dev/null +++ b/Patatrack/Tutorial/test/benchmarkCartesianVectors.py @@ -0,0 +1,23 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("PRINT") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 1 ), + numberOfStreams = cms.untracked.uint32( 1 ), + wantSummary = cms.untracked.bool( True ) +) + +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring("file:cylindricalVectors.root"), +) + +process.convertToCartesianVectors = cms.EDProducer('ConvertToCartesianVectors', + input = cms.InputTag('generateCylindricalVectors') +) + +process.path = cms.Path(process.convertToCartesianVectors) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( -1 ) +) diff --git a/Patatrack/Tutorial/test/benchmarkCartesianVectorsCUDA.py b/Patatrack/Tutorial/test/benchmarkCartesianVectorsCUDA.py new file mode 100644 index 0000000000000..4a569be234bd2 --- /dev/null +++ b/Patatrack/Tutorial/test/benchmarkCartesianVectorsCUDA.py @@ -0,0 +1,27 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("PRINT") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 1 ), + numberOfStreams = cms.untracked.uint32( 1 ), + wantSummary = cms.untracked.bool( True ) +) + +process.load("FWCore.MessageService.MessageLogger_cfi") +process.load("HeterogeneousCore.CUDAServices.CUDAService_cfi") +process.MessageLogger.categories.append("CUDAService") + +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring("file:cylindricalVectors.root"), +) + +process.convertToCartesianVectors = cms.EDProducer('ConvertToCartesianVectorsCUDA', + input = cms.InputTag('generateCylindricalVectors') +) + +process.path = cms.Path(process.convertToCartesianVectors) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( -1 ) +) diff --git a/Patatrack/Tutorial/test/compareCartesianVectors.py b/Patatrack/Tutorial/test/compareCartesianVectors.py new file mode 100644 index 0000000000000..2df601440e309 --- /dev/null +++ b/Patatrack/Tutorial/test/compareCartesianVectors.py @@ -0,0 +1,37 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("PRINT") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 1 ), + numberOfStreams = cms.untracked.uint32( 1 ), + wantSummary = cms.untracked.bool( False ) +) + +process.load("FWCore.MessageService.MessageLogger_cfi") +process.load("HeterogeneousCore.CUDAServices.CUDAService_cfi") +process.MessageLogger.categories.append("CUDAService") + +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring("file:cylindricalVectors.root"), +) + +process.convertToCartesianVectorsCUDA = cms.EDProducer('ConvertToCartesianVectorsCUDA', + input = cms.InputTag('generateCylindricalVectors') +) + +process.convertToCartesianVectors = cms.EDProducer('ConvertToCartesianVectors', + input = cms.InputTag('generateCylindricalVectors') +) + +process.compareCartesianVectors = cms.EDAnalyzer('CompareCartesianVectors', + first = cms.InputTag('convertToCartesianVectorsCUDA'), + second = cms.InputTag('convertToCartesianVectors'), + precision = cms.double(1.e-7) +) + +process.path = cms.Path(process.convertToCartesianVectorsCUDA + process.convertToCartesianVectors + process.compareCartesianVectors) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( -1 ) +) diff --git a/Patatrack/Tutorial/test/generateCylindricalVectors.py b/Patatrack/Tutorial/test/generateCylindricalVectors.py new file mode 100644 index 0000000000000..854a57803323f --- /dev/null +++ b/Patatrack/Tutorial/test/generateCylindricalVectors.py @@ -0,0 +1,30 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("GEN") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 8 ), + numberOfStreams = cms.untracked.uint32( 0 ), + wantSummary = cms.untracked.bool( True ) +) + +process.source = cms.Source("EmptySource") + +process.generateCylindricalVectors = cms.EDProducer('GenerateCylindricalVectors', + size = cms.uint32(10000) +) + +process.path = cms.Path(process.generateCylindricalVectors) + +process.out = cms.OutputModule("PoolOutputModule", + fileName = cms.untracked.string("cylindricalVectors.root"), + outputCommands = cms.untracked.vstring( + 'drop *', + 'keep *_generateCylindricalVectors_*_*') +) + +process.endp = cms.EndPath(process.out) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( 1200 ) +) diff --git a/Patatrack/Tutorial/test/printCartesianVectors.py b/Patatrack/Tutorial/test/printCartesianVectors.py new file mode 100644 index 0000000000000..00f3a8f1c35e4 --- /dev/null +++ b/Patatrack/Tutorial/test/printCartesianVectors.py @@ -0,0 +1,27 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("PRINT") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 1 ), + numberOfStreams = cms.untracked.uint32( 1 ), + wantSummary = cms.untracked.bool( False ) +) + +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring("file:cylindricalVectors.root"), +) + +process.convertToCartesianVectors = cms.EDProducer('ConvertToCartesianVectors', + input = cms.InputTag('generateCylindricalVectors') +) + +process.printCartesianVectors = cms.EDAnalyzer('PrintCartesianVectors', + input = cms.InputTag('convertToCartesianVectors') +) + +process.path = cms.Path(process.convertToCartesianVectors + process.printCartesianVectors) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( 1 ) +) diff --git a/Patatrack/Tutorial/test/printCartesianVectorsCUDA.py b/Patatrack/Tutorial/test/printCartesianVectorsCUDA.py new file mode 100644 index 0000000000000..78ef4d05c56b1 --- /dev/null +++ b/Patatrack/Tutorial/test/printCartesianVectorsCUDA.py @@ -0,0 +1,31 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("PRINT") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 1 ), + numberOfStreams = cms.untracked.uint32( 1 ), + wantSummary = cms.untracked.bool( False ) +) + +process.load("FWCore.MessageService.MessageLogger_cfi") +process.load("HeterogeneousCore.CUDAServices.CUDAService_cfi") +process.MessageLogger.categories.append("CUDAService") + +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring("file:cylindricalVectors.root"), +) + +process.convertToCartesianVectors = cms.EDProducer('ConvertToCartesianVectorsCUDA', + input = cms.InputTag('generateCylindricalVectors') +) + +process.printCartesianVectors = cms.EDAnalyzer('PrintCartesianVectors', + input = cms.InputTag('convertToCartesianVectors') +) + +process.path = cms.Path(process.convertToCartesianVectors + process.printCartesianVectors) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( 1 ) +) diff --git a/Patatrack/Tutorial/test/printCylindricalVectors.py b/Patatrack/Tutorial/test/printCylindricalVectors.py new file mode 100644 index 0000000000000..d74db04a367dd --- /dev/null +++ b/Patatrack/Tutorial/test/printCylindricalVectors.py @@ -0,0 +1,23 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("PRINT") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 1 ), + numberOfStreams = cms.untracked.uint32( 1 ), + wantSummary = cms.untracked.bool( False ) +) + +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring("file:cylindricalVectors.root"), +) + +process.printCylindricalVectors = cms.EDAnalyzer('PrintCylindricalVectors', + input = cms.InputTag('generateCylindricalVectors') +) + +process.path = cms.Path(process.printCylindricalVectors) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( 1 ) +) From ae56a932d13ec9c8ee79446c27e791bd6e3a4087 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 11 Dec 2019 14:49:09 +0100 Subject: [PATCH 3/7] Solution to the first exercise --- .../plugins/ConvertToCartesianVectorsCUDA.cc | 14 +++++++++----- Patatrack/Tutorial/plugins/cudavectors.cu | 16 ++++++++++++---- 2 files changed, 21 insertions(+), 9 deletions(-) diff --git a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc index 598a10a69a32e..4877b2a782a0a 100644 --- a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc +++ b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc @@ -43,19 +43,23 @@ void ConvertToCartesianVectorsCUDA::produce(edm::Event& event, const edm::EventS auto product = std::make_unique(elements); // allocate memory on the GPU for the cylindrical and cartesian vectors - // fill here ... + cudavectors::CylindricalVector* gpu_input; + cudavectors::CartesianVector* gpu_product; + cudaCheck(cudaMalloc(&gpu_input, sizeof(cudavectors::CylindricalVector) * elements)); + cudaCheck(cudaMalloc(&gpu_product, sizeof(cudavectors::CartesianVector) * elements)); // copy the input data to the GPU - // fill here ... + cudaCheck(cudaMemcpy(gpu_input, input.data(), sizeof(cudavectors::CylindricalVector) * elements, cudaMemcpyHostToDevice)); // convert the vectors from cylindrical to cartesian coordinates, on the GPU - // fill here ... + cudavectors::convertWrapper(gpu_input, gpu_product, elements); // copy the result from the GPU - // fill here ... + cudaCheck(cudaMemcpy(product->data(), gpu_product, sizeof(cudavectors::CartesianVector) * elements, cudaMemcpyDeviceToHost)); // free the GPU memory - // fill here ... + cudaCheck(cudaFree(gpu_input)); + cudaCheck(cudaFree(gpu_product)); event.put(output_, std::move(product)); } diff --git a/Patatrack/Tutorial/plugins/cudavectors.cu b/Patatrack/Tutorial/plugins/cudavectors.cu index 8da23582aad2c..53a5418320a4b 100644 --- a/Patatrack/Tutorial/plugins/cudavectors.cu +++ b/Patatrack/Tutorial/plugins/cudavectors.cu @@ -11,16 +11,24 @@ namespace cudavectors { __host__ __device__ inline void convert(CylindricalVector const& cylindrical, CartesianVector & cartesian) { - // fill here ... + cartesian.x = cylindrical.rho * std::cos(cylindrical.phi); + cartesian.y = cylindrical.rho * std::sin(cylindrical.phi); + cartesian.z = cylindrical.rho * std::sinh(cylindrical.eta); } __global__ void convertKernel(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size) { - // fill here ... + auto firstElement = threadIdx.x + blockIdx.x * blockDim.x; + auto gridSize = blockDim.x * gridDim.x; + + for (size_t i = firstElement; i < size; i += gridSize) { + convert(cylindrical[i], cartesian[i]); + } } void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size) { - // fill here ... - //convertKernel<<>>(cylindrical, cartesian, size); + auto blockSize = 512; // somewhat arbitrary + auto gridSize = (size + blockSize - 1) / blockSize; // round up to cover the sample size + convertKernel<<>>(cylindrical, cartesian, size); cudaCheck(cudaGetLastError()); } From 86aa6405808b9a6a832c19d604cda6e28d0f81b9 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 11 Dec 2019 14:58:28 +0100 Subject: [PATCH 4/7] Solution to the second exercise --- Patatrack/Tutorial/plugins/BuildFile.xml | 1 + .../plugins/ConvertToCartesianVectorsCUDA.cc | 22 +++++++++---------- 2 files changed, 12 insertions(+), 11 deletions(-) diff --git a/Patatrack/Tutorial/plugins/BuildFile.xml b/Patatrack/Tutorial/plugins/BuildFile.xml index 5e307e5c27188..78b20778a3e8a 100644 --- a/Patatrack/Tutorial/plugins/BuildFile.xml +++ b/Patatrack/Tutorial/plugins/BuildFile.xml @@ -3,4 +3,5 @@ + diff --git a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc index 4877b2a782a0a..7bc4537b58c21 100644 --- a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc +++ b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc @@ -12,6 +12,8 @@ #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "FWCore/Utilities/interface/StreamID.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" +#include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" #include "cudavectors.h" @@ -43,23 +45,21 @@ void ConvertToCartesianVectorsCUDA::produce(edm::Event& event, const edm::EventS auto product = std::make_unique(elements); // allocate memory on the GPU for the cylindrical and cartesian vectors - cudavectors::CylindricalVector* gpu_input; - cudavectors::CartesianVector* gpu_product; - cudaCheck(cudaMalloc(&gpu_input, sizeof(cudavectors::CylindricalVector) * elements)); - cudaCheck(cudaMalloc(&gpu_product, sizeof(cudavectors::CartesianVector) * elements)); + auto gpu_input = cudautils::make_device_unique(elements, cudaStreamDefault); + auto gpu_product = cudautils::make_device_unique(elements, cudaStreamDefault); + + // allocate memory on the CPU for the transfer buffer + auto cpu_input = cudautils::make_host_noncached_unique(elements, cudaHostAllocWriteCombined); + std::memcpy(cpu_input.get(), input.data(), sizeof(cudavectors::CylindricalVector) * elements); // copy the input data to the GPU - cudaCheck(cudaMemcpy(gpu_input, input.data(), sizeof(cudavectors::CylindricalVector) * elements, cudaMemcpyHostToDevice)); + cudaCheck(cudaMemcpy(gpu_input.get(), cpu_input.get(), sizeof(cudavectors::CylindricalVector) * elements, cudaMemcpyHostToDevice)); // convert the vectors from cylindrical to cartesian coordinates, on the GPU - cudavectors::convertWrapper(gpu_input, gpu_product, elements); + cudavectors::convertWrapper(gpu_input.get(), gpu_product.get(), elements); // copy the result from the GPU - cudaCheck(cudaMemcpy(product->data(), gpu_product, sizeof(cudavectors::CartesianVector) * elements, cudaMemcpyDeviceToHost)); - - // free the GPU memory - cudaCheck(cudaFree(gpu_input)); - cudaCheck(cudaFree(gpu_product)); + cudaCheck(cudaMemcpy(product->data(), gpu_product.get(), sizeof(cudavectors::CartesianVector) * elements, cudaMemcpyDeviceToHost)); event.put(output_, std::move(product)); } From b39c5ba91d62cf02ae81231c67b3480bc370521e Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 12 Dec 2019 11:03:50 +0100 Subject: [PATCH 5/7] Solution to the third exercise --- Patatrack/Tutorial/plugins/BuildFile.xml | 1 + .../plugins/ConvertToCartesianVectorsCUDA.cc | 42 ++++++++++++++----- Patatrack/Tutorial/plugins/cudavectors.cu | 4 +- Patatrack/Tutorial/plugins/cudavectors.h | 2 +- 4 files changed, 36 insertions(+), 13 deletions(-) diff --git a/Patatrack/Tutorial/plugins/BuildFile.xml b/Patatrack/Tutorial/plugins/BuildFile.xml index 78b20778a3e8a..218c1b4791438 100644 --- a/Patatrack/Tutorial/plugins/BuildFile.xml +++ b/Patatrack/Tutorial/plugins/BuildFile.xml @@ -3,5 +3,6 @@ + diff --git a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc index 7bc4537b58c21..c47277f117768 100644 --- a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc +++ b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc @@ -11,13 +11,15 @@ #include "FWCore/Framework/interface/stream/EDProducer.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "FWCore/Utilities/interface/StreamID.h" +#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" +#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" #include "cudavectors.h" -class ConvertToCartesianVectorsCUDA : public edm::stream::EDProducer<> { +class ConvertToCartesianVectorsCUDA : public edm::stream::EDProducer { public: explicit ConvertToCartesianVectorsCUDA(const edm::ParameterSet&); ~ConvertToCartesianVectorsCUDA() = default; @@ -28,10 +30,12 @@ class ConvertToCartesianVectorsCUDA : public edm::stream::EDProducer<> { using CartesianVectors = std::vector; using CylindricalVectors = std::vector; - virtual void produce(edm::Event&, const edm::EventSetup&) override; + void acquire(edm::Event const& event, edm::EventSetup const& setup, edm::WaitingTaskWithArenaHolder waitingTaskHolder) override; + void produce(edm::Event& event, edm::EventSetup const& setup) override; edm::EDGetTokenT input_; edm::EDPutTokenT output_; + cudautils::host::unique_ptr output_buffer_; }; ConvertToCartesianVectorsCUDA::ConvertToCartesianVectorsCUDA(const edm::ParameterSet& config) @@ -39,28 +43,46 @@ ConvertToCartesianVectorsCUDA::ConvertToCartesianVectorsCUDA(const edm::Paramete output_ = produces(); } -void ConvertToCartesianVectorsCUDA::produce(edm::Event& event, const edm::EventSetup& setup) { +void ConvertToCartesianVectorsCUDA::acquire(const edm::Event& event, const edm::EventSetup& setup, edm::WaitingTaskWithArenaHolder waitingTaskHolder) { + // set the current device and create a CUDA stream + CUDAScopedContextAcquire ctx{event.streamID(), std::move(waitingTaskHolder)}; + auto const& input = event.get(input_); auto elements = input.size(); - auto product = std::make_unique(elements); // allocate memory on the GPU for the cylindrical and cartesian vectors - auto gpu_input = cudautils::make_device_unique(elements, cudaStreamDefault); - auto gpu_product = cudautils::make_device_unique(elements, cudaStreamDefault); + auto gpu_input = cudautils::make_device_unique(elements, ctx.stream()); + auto gpu_product = cudautils::make_device_unique(elements, ctx.stream()); - // allocate memory on the CPU for the transfer buffer + // allocate memory on the CPU for the transfer buffers auto cpu_input = cudautils::make_host_noncached_unique(elements, cudaHostAllocWriteCombined); + output_buffer_ = cudautils::make_host_unique(elements, ctx.stream()); std::memcpy(cpu_input.get(), input.data(), sizeof(cudavectors::CylindricalVector) * elements); // copy the input data to the GPU - cudaCheck(cudaMemcpy(gpu_input.get(), cpu_input.get(), sizeof(cudavectors::CylindricalVector) * elements, cudaMemcpyHostToDevice)); + cudaCheck(cudaMemcpyAsync(gpu_input.get(), cpu_input.get(), sizeof(cudavectors::CylindricalVector) * elements, cudaMemcpyHostToDevice, ctx.stream())); // convert the vectors from cylindrical to cartesian coordinates, on the GPU - cudavectors::convertWrapper(gpu_input.get(), gpu_product.get(), elements); + cudavectors::convertWrapper(gpu_input.get(), gpu_product.get(), elements, ctx.stream()); // copy the result from the GPU - cudaCheck(cudaMemcpy(product->data(), gpu_product.get(), sizeof(cudavectors::CartesianVector) * elements, cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpyAsync(output_buffer_.get(), gpu_product.get(), sizeof(cudavectors::CartesianVector) * elements, cudaMemcpyDeviceToHost, ctx.stream())); + + // the CUDA context automatically sets up a callback to notify the framework when the operations on the CUDA stream are complete +} + +void ConvertToCartesianVectorsCUDA::produce(edm::Event& event, const edm::EventSetup& setup) { + // no need for a CUDA context here, because there are no CUDA operations + + auto const& input = event.get(input_); + auto elements = input.size(); + + // instantiate the event product, copy the results from the output buffer, and free it + auto product = std::make_unique(elements); + std::memcpy((void*) product->data(), output_buffer_.get(), sizeof(cudavectors::CartesianVector) * elements); + output_buffer_.reset(); + // put the product in the event event.put(output_, std::move(product)); } diff --git a/Patatrack/Tutorial/plugins/cudavectors.cu b/Patatrack/Tutorial/plugins/cudavectors.cu index 53a5418320a4b..102783e0a7704 100644 --- a/Patatrack/Tutorial/plugins/cudavectors.cu +++ b/Patatrack/Tutorial/plugins/cudavectors.cu @@ -25,10 +25,10 @@ namespace cudavectors { } } - void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size) { + void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size, cudaStream_t stream) { auto blockSize = 512; // somewhat arbitrary auto gridSize = (size + blockSize - 1) / blockSize; // round up to cover the sample size - convertKernel<<>>(cylindrical, cartesian, size); + convertKernel<<>>(cylindrical, cartesian, size); cudaCheck(cudaGetLastError()); } diff --git a/Patatrack/Tutorial/plugins/cudavectors.h b/Patatrack/Tutorial/plugins/cudavectors.h index c20625f37b2b0..5d309e4a3f772 100644 --- a/Patatrack/Tutorial/plugins/cudavectors.h +++ b/Patatrack/Tutorial/plugins/cudavectors.h @@ -15,7 +15,7 @@ namespace cudavectors { float z; }; - void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size); + void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size, cudaStream_t stream); } // namespace cudavectors From a9295025e378a71b2621b47755cad4eacc81f788 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 11 Dec 2019 13:51:01 +0100 Subject: [PATCH 6/7] Solution to the fourth exercise --- .../test/printCartesianVectorsAuto.py | 37 +++++++++++++++++++ 1 file changed, 37 insertions(+) create mode 100644 Patatrack/Tutorial/test/printCartesianVectorsAuto.py diff --git a/Patatrack/Tutorial/test/printCartesianVectorsAuto.py b/Patatrack/Tutorial/test/printCartesianVectorsAuto.py new file mode 100644 index 0000000000000..7781a7d4a7bd6 --- /dev/null +++ b/Patatrack/Tutorial/test/printCartesianVectorsAuto.py @@ -0,0 +1,37 @@ +import FWCore.ParameterSet.Config as cms + +process = cms.Process("PRINT") + +process.options = cms.untracked.PSet( + numberOfThreads = cms.untracked.uint32( 1 ), + numberOfStreams = cms.untracked.uint32( 1 ), + wantSummary = cms.untracked.bool( False ) +) + +process.load("FWCore.MessageService.MessageLogger_cfi") +process.load("HeterogeneousCore.CUDAServices.CUDAService_cfi") +process.MessageLogger.categories.append("CUDAService") + +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring("file:cylindricalVectors.root"), +) + +from HeterogeneousCore.CUDACore.SwitchProducerCUDA import SwitchProducerCUDA +process.convertToCartesianVectors = SwitchProducerCUDA( + cpu = cms.EDProducer('ConvertToCartesianVectors', + input = cms.InputTag('generateCylindricalVectors') + ), + cuda = cms.EDProducer('ConvertToCartesianVectorsCUDA', + input = cms.InputTag('generateCylindricalVectors') + ) +) + +process.printCartesianVectors = cms.EDAnalyzer('PrintCartesianVectors', + input = cms.InputTag('convertToCartesianVectors') +) + +process.path = cms.Path(process.convertToCartesianVectors + process.printCartesianVectors) + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32( 1 ) +) From d477b773fe26504e923bbd8427c53e628114711c Mon Sep 17 00:00:00 2001 From: aravind sugunan Date: Tue, 23 Jun 2020 09:03:02 +0200 Subject: [PATCH 7/7] adding updated namspace info for the solution in 'https://patatrack.web.cern.ch/patatrack/wiki/cuda_training_dpg_12_2019_part2/' --- DataFormats/Math/src/classes_def.xml | 2 + Patatrack/Tutorial/plugins/BuildFile.xml | 2 +- .../plugins/ConvertToCartesianVectorsCUDA.cc | 74 +++++++++---------- Patatrack/Tutorial/plugins/cudavectors.cu | 13 ++-- Patatrack/Tutorial/plugins/cudavectors.h | 2 +- 5 files changed, 49 insertions(+), 44 deletions(-) diff --git a/DataFormats/Math/src/classes_def.xml b/DataFormats/Math/src/classes_def.xml index 51b3abbab4790..36213287e4bb3 100755 --- a/DataFormats/Math/src/classes_def.xml +++ b/DataFormats/Math/src/classes_def.xml @@ -4,6 +4,8 @@ + + diff --git a/Patatrack/Tutorial/plugins/BuildFile.xml b/Patatrack/Tutorial/plugins/BuildFile.xml index 218c1b4791438..b6607d1c361a9 100644 --- a/Patatrack/Tutorial/plugins/BuildFile.xml +++ b/Patatrack/Tutorial/plugins/BuildFile.xml @@ -3,6 +3,6 @@ - + diff --git a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc index c47277f117768..2e88c02e762cb 100644 --- a/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc +++ b/Patatrack/Tutorial/plugins/ConvertToCartesianVectorsCUDA.cc @@ -11,12 +11,12 @@ #include "FWCore/Framework/interface/stream/EDProducer.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "FWCore/Utilities/interface/StreamID.h" -#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" + +#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" -#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" - +#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" #include "cudavectors.h" class ConvertToCartesianVectorsCUDA : public edm::stream::EDProducer { @@ -31,11 +31,10 @@ class ConvertToCartesianVectorsCUDA : public edm::stream::EDProducer; void acquire(edm::Event const& event, edm::EventSetup const& setup, edm::WaitingTaskWithArenaHolder waitingTaskHolder) override; - void produce(edm::Event& event, edm::EventSetup const& setup) override; - + virtual void produce(edm::Event&, const edm::EventSetup&) override; + cms::cuda::host::unique_ptr output_buffer_; edm::EDGetTokenT input_; edm::EDPutTokenT output_; - cudautils::host::unique_ptr output_buffer_; }; ConvertToCartesianVectorsCUDA::ConvertToCartesianVectorsCUDA(const edm::ParameterSet& config) @@ -43,47 +42,49 @@ ConvertToCartesianVectorsCUDA::ConvertToCartesianVectorsCUDA(const edm::Paramete output_ = produces(); } -void ConvertToCartesianVectorsCUDA::acquire(const edm::Event& event, const edm::EventSetup& setup, edm::WaitingTaskWithArenaHolder waitingTaskHolder) { - // set the current device and create a CUDA stream - CUDAScopedContextAcquire ctx{event.streamID(), std::move(waitingTaskHolder)}; + +void ConvertToCartesianVectorsCUDA::acquire(const edm::Event& event, const edm::EventSetup& setup, + edm::WaitingTaskWithArenaHolder waitingTaskHolder) { + + cms::cuda::ScopedContextAcquire ctx{event.streamID(), std::move(waitingTaskHolder)}; auto const& input = event.get(input_); auto elements = input.size(); // allocate memory on the GPU for the cylindrical and cartesian vectors - auto gpu_input = cudautils::make_device_unique(elements, ctx.stream()); - auto gpu_product = cudautils::make_device_unique(elements, ctx.stream()); - - // allocate memory on the CPU for the transfer buffers - auto cpu_input = cudautils::make_host_noncached_unique(elements, cudaHostAllocWriteCombined); - output_buffer_ = cudautils::make_host_unique(elements, ctx.stream()); - std::memcpy(cpu_input.get(), input.data(), sizeof(cudavectors::CylindricalVector) * elements); - + + auto gpu_input = cms::cuda::make_device_unique(elements, ctx.stream()); + auto gpu_product = cms::cuda::make_device_unique(elements, ctx.stream()); + auto cpu_input = cms::cuda::make_host_noncached_unique(elements, cudaHostAllocWriteCombined); + output_buffer_ = cms::cuda::make_host_unique(elements, ctx.stream()); + // copy the input data to the GPU + + std::memcpy(cpu_input.get(), input.data(), sizeof(cudavectors::CylindricalVector) * elements); cudaCheck(cudaMemcpyAsync(gpu_input.get(), cpu_input.get(), sizeof(cudavectors::CylindricalVector) * elements, cudaMemcpyHostToDevice, ctx.stream())); - // convert the vectors from cylindrical to cartesian coordinates, on the GPU - cudavectors::convertWrapper(gpu_input.get(), gpu_product.get(), elements, ctx.stream()); - + // convert the vectors from cylindrical to cartesian coordinates, on the GPU + + cudavectors::convertWrapper(gpu_input.get(), gpu_product.get(), elements, ctx.stream()); + // copy the result from the GPU - cudaCheck(cudaMemcpyAsync(output_buffer_.get(), gpu_product.get(), sizeof(cudavectors::CartesianVector) * elements, cudaMemcpyDeviceToHost, ctx.stream())); - - // the CUDA context automatically sets up a callback to notify the framework when the operations on the CUDA stream are complete -} - + + cudaCheck(cudaMemcpyAsync(output_buffer_.get(), gpu_product.get(), sizeof(cudavectors::CartesianVector) * elements, cudaMemcpyDeviceToHost, ctx.stream())); + + // free the GPU memory + // no need of explicit free operation +} void ConvertToCartesianVectorsCUDA::produce(edm::Event& event, const edm::EventSetup& setup) { - // no need for a CUDA context here, because there are no CUDA operations - - auto const& input = event.get(input_); - auto elements = input.size(); - + //no need for a CUDA context here, because there are no CUDA operations + + auto const& input = event.get(input_); + auto elements = input.size(); + // instantiate the event product, copy the results from the output buffer, and free it - auto product = std::make_unique(elements); - std::memcpy((void*) product->data(), output_buffer_.get(), sizeof(cudavectors::CartesianVector) * elements); - output_buffer_.reset(); - - // put the product in the event - event.put(output_, std::move(product)); + auto product = std::make_unique(elements); + std::memcpy((void*) product->data(), output_buffer_.get(), sizeof(cudavectors::CartesianVector) * elements); + output_buffer_.reset(); + event.put(output_, std::move(product)); } void ConvertToCartesianVectorsCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { @@ -91,6 +92,5 @@ void ConvertToCartesianVectorsCUDA::fillDescriptions(edm::ConfigurationDescripti desc.add("input", edm::InputTag("cylindricalVectors")); descriptions.addWithDefaultLabel(desc); } - // define this as a plug-in DEFINE_FWK_MODULE(ConvertToCartesianVectorsCUDA); diff --git a/Patatrack/Tutorial/plugins/cudavectors.cu b/Patatrack/Tutorial/plugins/cudavectors.cu index 102783e0a7704..490a315ee4ff6 100644 --- a/Patatrack/Tutorial/plugins/cudavectors.cu +++ b/Patatrack/Tutorial/plugins/cudavectors.cu @@ -11,21 +11,24 @@ namespace cudavectors { __host__ __device__ inline void convert(CylindricalVector const& cylindrical, CartesianVector & cartesian) { + cartesian.x = cylindrical.rho * std::cos(cylindrical.phi); cartesian.y = cylindrical.rho * std::sin(cylindrical.phi); - cartesian.z = cylindrical.rho * std::sinh(cylindrical.eta); - } + cartesian.z = cylindrical.rho * std::sinh(cylindrical.eta); + + } __global__ void convertKernel(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size) { + auto firstElement = threadIdx.x + blockIdx.x * blockDim.x; auto gridSize = blockDim.x * gridDim.x; for (size_t i = firstElement; i < size; i += gridSize) { convert(cylindrical[i], cartesian[i]); } - } - - void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size, cudaStream_t stream) { + } + void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size,cudaStream_t stream) { + //convertKernel<<>>(cylindrical, cartesian, size); auto blockSize = 512; // somewhat arbitrary auto gridSize = (size + blockSize - 1) / blockSize; // round up to cover the sample size convertKernel<<>>(cylindrical, cartesian, size); diff --git a/Patatrack/Tutorial/plugins/cudavectors.h b/Patatrack/Tutorial/plugins/cudavectors.h index 5d309e4a3f772..7aac345e86ffe 100644 --- a/Patatrack/Tutorial/plugins/cudavectors.h +++ b/Patatrack/Tutorial/plugins/cudavectors.h @@ -15,7 +15,7 @@ namespace cudavectors { float z; }; - void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size, cudaStream_t stream); + void convertWrapper(CylindricalVector const* cylindrical, CartesianVector* cartesian, size_t size,cudaStream_t stream); } // namespace cudavectors