Inheritance diagram for TestCUDAProducerGPUEWTask:

Public Member Functions
void	acquire (edm::Event const &iEvent, edm::EventSetup const &iSetup, edm::WaitingTaskWithArenaHolder waitingTaskHolder) override

void	produce (edm::Event &iEvent, const edm::EventSetup &iSetup) override

	TestCUDAProducerGPUEWTask (edm::ParameterSet const &iConfig)

	~TestCUDAProducerGPUEWTask () override=default

Public Member Functions inherited from edm::stream::EDProducer< edm::ExternalWork >
	EDProducer ()=default

	EDProducer (const EDProducer &)=delete

bool	hasAbilityToProduceInBeginLumis () const final

bool	hasAbilityToProduceInBeginProcessBlocks () const final

bool	hasAbilityToProduceInBeginRuns () const final

bool	hasAbilityToProduceInEndLumis () const final

bool	hasAbilityToProduceInEndProcessBlocks () const final

bool	hasAbilityToProduceInEndRuns () const final

const EDProducer &	operator= (const EDProducer &)=delete

Static Public Member Functions
static void	fillDescriptions (edm::ConfigurationDescriptions &descriptions)

Private Member Functions
void	addSimpleWork (edm::EventNumber_t eventID, edm::StreamID streamID, cms::cuda::ScopedContextTask &ctx)

Private Attributes
cms::cuda::ContextState	ctxState_

cms::cuda::device::unique_ptr < float[]>	devicePtr_

edm::EDPutTokenT < cms::cuda::Product < cms::cudatest::Thing > > const	dstToken_

TestCUDAProducerGPUKernel	gpuAlgo_

cms::cuda::host::noncached::unique_ptr < float >	hostData_

std::string const	label_

edm::EDGetTokenT < cms::cuda::Product < cms::cudatest::Thing > > const	srcToken_

Additional Inherited Members
Public Types inherited from edm::stream::EDProducer< edm::ExternalWork >
using	CacheTypes = CacheContexts< T...>

using	GlobalCache = typename CacheTypes::GlobalCache

using	HasAbility = AbilityChecker< T...>

using	InputProcessBlockCache = typename CacheTypes::InputProcessBlockCache

using	LuminosityBlockCache = typename CacheTypes::LuminosityBlockCache

using	LuminosityBlockContext = LuminosityBlockContextT< LuminosityBlockCache, RunCache, GlobalCache >

using	LuminosityBlockSummaryCache = typename CacheTypes::LuminosityBlockSummaryCache

using	RunCache = typename CacheTypes::RunCache

using	RunContext = RunContextT< RunCache, GlobalCache >

using	RunSummaryCache = typename CacheTypes::RunSummaryCache

Detailed Description

Definition at line 24 of file TestCUDAProducerGPUEWTask.cc.

Constructor & Destructor Documentation

TestCUDAProducerGPUEWTask::TestCUDAProducerGPUEWTask ( edm::ParameterSet const & iConfig )

explicit

Definition at line 48 of file TestCUDAProducerGPUEWTask.cc.

References fwrapper::cs, dstToken_, CUDAService::enabled(), edm::ParameterSet::getParameter(), hostData_, srcToken_, and AlCaHLTBitMon_QueryRunRegistry::string.

     : label_{iConfig.getParameter<std::string>("@module_label")},
       srcToken_{consumes<cms::cuda::Product<cms::cudatest::Thing>>(iConfig.getParameter<edm::InputTag>("src"))},
       dstToken_{produces<cms::cuda::Product<cms::cudatest::Thing>>()} {
   edm::Service<CUDAService> cs;
   if (cs->enabled()) {
     hostData_ = cms::cuda::make_host_noncached_unique<float>();
   }
 }

TestCUDAProducerGPUEWTask::~TestCUDAProducerGPUEWTask ( )

overridedefault

Member Function Documentation

void TestCUDAProducerGPUEWTask::acquire	(	edm::Event const &	iEvent,
		edm::EventSetup const &	iSetup,
		edm::WaitingTaskWithArenaHolder	waitingTaskHolder
	)

override

Definition at line 71 of file TestCUDAProducerGPUEWTask.cc.

References addSimpleWork(), ctxState_, cudaCheck, devicePtr_, edm::EventID::event(), cms::cudatest::Thing::get(), edm::Event::get(), gpuAlgo_, hostData_, edm::EventBase::id(), iev, recoMuon::in, input, label_, TestCUDAProducerGPUKernel::runAlgo(), srcToken_, and edm::Event::streamID().

                                                                                          {
   edm::LogVerbatim("TestCUDAProducerGPUEWTask") << label_ << " TestCUDAProducerGPUEWTask::acquire begin event "
                                                 << iEvent.id().event() << " stream " << iEvent.streamID();
 
   auto const& in = iEvent.get(srcToken_);
   cms::cuda::ScopedContextAcquire ctx{in, waitingTaskHolder, ctxState_};
 
   cms::cudatest::Thing const& input = ctx.get(in);
 
   devicePtr_ = gpuAlgo_.runAlgo(label_, input.get(), ctx.stream());
   // Mimick the need to transfer some of the GPU data back to CPU to
   // be used for something within this module, or to be put in the
   // event.
   cudaCheck(
       cudaMemcpyAsync(hostData_.get(), devicePtr_.get() + 10, sizeof(float), cudaMemcpyDeviceToHost, ctx.stream()));
   // Push a task to run addSimpleWork() after the asynchronous work
   // (and acquire()) has finished instead of produce()
   ctx.pushNextTask([iev = iEvent.id().event(), istr = iEvent.streamID(), this](cms::cuda::ScopedContextTask ctx) {
     addSimpleWork(iev, istr, ctx);
   });
 
   edm::LogVerbatim("TestCUDAProducerGPUEWTask") << label_ << " TestCUDAProducerGPUEWTask::acquire end event "
                                                 << iEvent.id().event() << " stream " << iEvent.streamID();
 }

void TestCUDAProducerGPUEWTask::addSimpleWork	(	edm::EventNumber_t	eventID,
		edm::StreamID	streamID,
		cms::cuda::ScopedContextTask &	ctx
	)

private

Definition at line 98 of file TestCUDAProducerGPUEWTask.cc.

References cudaCheck, devicePtr_, gpuAlgo_, hostData_, label_, cms::cuda::ScopedContextTask::pushNextTask(), TestCUDAProducerGPUKernel::runSimpleAlgo(), and cms::cuda::impl::ScopedContextBase::stream().

Referenced by acquire().

                                                                              {
   if (*hostData_ < 13) {
     edm::LogVerbatim("TestCUDAProducerGPUEWTask")
         << label_ << " TestCUDAProducerGPUEWTask::addSimpleWork begin event " << eventID << " stream " << streamID
         << " 10th element " << *hostData_ << " not satisfied, queueing more work";
     cudaCheck(
         cudaMemcpyAsync(hostData_.get(), devicePtr_.get() + 10, sizeof(float), cudaMemcpyDeviceToHost, ctx.stream()));
 
     ctx.pushNextTask(
         [eventID, streamID, this](cms::cuda::ScopedContextTask ctx) { addSimpleWork(eventID, streamID, ctx); });
     gpuAlgo_.runSimpleAlgo(devicePtr_.get(), ctx.stream());
     edm::LogVerbatim("TestCUDAProducerGPUEWTask")
         << label_ << " TestCUDAProducerGPUEWTask::addSimpleWork end event " << eventID << " stream " << streamID;
   } else {
     edm::LogVerbatim("TestCUDAProducerGPUEWTask")
         << label_ << " TestCUDAProducerGPUEWTask::addSimpleWork event " << eventID << " stream " << streamID
         << " 10th element " << *hostData_ << " not queueing more work";
   }
 }

void TestCUDAProducerGPUEWTask::fillDescriptions ( edm::ConfigurationDescriptions & descriptions )

static

Definition at line 58 of file TestCUDAProducerGPUEWTask.cc.

References edm::ParameterSetDescription::add(), edm::ConfigurationDescriptions::addWithDefaultLabel(), submitPVResolutionJobs::desc, HLT_FULL_cff::InputTag, and edm::ConfigurationDescriptions::setComment().

                                                                                            {
   edm::ParameterSetDescription desc;
   desc.add<edm::InputTag>("src", edm::InputTag());
   descriptions.addWithDefaultLabel(desc);
   descriptions.setComment(
       "This EDProducer is part of the TestCUDAProducer* family. It models a GPU algorithm this is not the first "
       "algorithm in the chain of the GPU EDProducers, and that transfers some data from GPU to CPU multiple times "
       "alternating the transfers and kernel executions (e.g. to decide which kernel to run next based on a value from "
       "GPU). A synchronization between GPU and CPU is needed after each transfer. The synchronizations are implemented "
       "with the ExternalWork extension and explicit TBB tasks within the module. Produces "
       "cms::cuda::Product<cms::cudatest::Thing>.");
 }

void TestCUDAProducerGPUEWTask::produce	(	edm::Event &	iEvent,
		const edm::EventSetup &	iSetup
	)

override

Definition at line 120 of file TestCUDAProducerGPUEWTask.cc.

References ctxState_, devicePtr_, dstToken_, edm::EventID::event(), Exception, hostData_, edm::EventBase::id(), label_, eostools::move(), and edm::Event::streamID().

                                                                                      {
   edm::LogVerbatim("TestCUDAProducerGPUEWTask")
       << label_ << " TestCUDAProducerGPUEWTask::produce begin event " << iEvent.id().event() << " stream "
       << iEvent.streamID() << " 10th element " << *hostData_;
   if (*hostData_ != 13) {
     throw cms::Exception("Assert") << "Expecting 10th element to be 13, got " << *hostData_;
   }
 
   cms::cuda::ScopedContextProduce ctx{ctxState_};
 
   ctx.emplace(iEvent, dstToken_, std::move(devicePtr_));
 
   edm::LogVerbatim("TestCUDAProducerGPUEWTask") << label_ << " TestCUDAProducerGPUEWTask::produce end event "
                                                 << iEvent.id().event() << " stream " << iEvent.streamID();
 }