Namespaces
	allocator

	detail

	device

	host

	impl

Classes
class	AtomicPairCounter

class	bad_alloc

class	ContextState

class	ESProduct

class	EventCache

class	FlexiStorage

class	FlexiStorage< I, -1 >

class	HistoContainer

class	HostAllocator

struct	LaunchParameters

class	LogAbsolute

class	LogError

class	LogImportant

class	LogInfo

class	LogPrint

class	LogProblem

class	LogSystem

class	LogVerbatim

class	LogWarning

class	MessageLogger

class	OneToManyAssoc

struct	OneToManyAssocView

class	PortableDeviceCollection

class	PortableHostCollection

class	Product

class	ProductBase

class	ScopedContextAcquire

class	ScopedContextAnalyze

class	ScopedContextProduce

class	ScopedContextTask

class	ScopedSetDevice

struct	SimpleVector

class	StreamCache

class	VecArray

Typedefs
using	SharedEventPtr = std::shared_ptr< std::remove_pointer_t< cudaEvent_t > >

using	SharedStreamPtr = std::shared_ptr< std::remove_pointer_t< cudaStream_t > >

Functions
template<typename Histo , typename T >
	__attribute__ ((always_inline)) void countFromVector(Histo *__restrict__ h

template<typename Assoc >
	__attribute__ ((always_inline)) void zeroAndInit(OneToManyAssocView< Assoc > view)

template<typename VT , typename T >
__host__ __device__	__attribute__ ((always_inline)) void blockPrefixScan(VT const *ci

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__	__attribute__ ((always_inline)) void forEachInBins(Hist const &hist

void	abortOnCudaError (const char file, int line, const char cmd, const char error, const char message, std::string_view description=std::string_view())

void	abortOnNvmlError (const char file, int line, const char cmd, const char error, const char message, std::string_view description=std::string_view())

void *	allocate_device (int dev, size_t nbytes, cudaStream_t stream)

void *	allocate_host (size_t nbytes, cudaStream_t stream)

	assert (be >=bs)

int	chooseDevice (edm::StreamID id)

template<typename T >
void	copyAsync (device::unique_ptr< T > &dst, const host::unique_ptr< T > &src, cudaStream_t stream)

template<typename T >
void	copyAsync (host::unique_ptr< T > &dst, const device::unique_ptr< T > &src, cudaStream_t stream)

template<typename T >
void	copyAsync (device::unique_ptr< T[]> &dst, const host::unique_ptr< T[]> &src, size_t nelements, cudaStream_t stream)

template<typename T >
void	copyAsync (host::unique_ptr< T[]> &dst, const device::unique_ptr< T[]> &src, size_t nelements, cudaStream_t stream)

template<typename T >
void	copyAsync (cms::cuda::device::unique_ptr< T[]> &dst, const std::vector< T, cms::cuda::HostAllocator< T >> &src, cudaStream_t stream)

template<typename T >
void	copyAsync (edm::propagate_const_array< cms::cuda::device::unique_ptr< T[]>> &dst, const std::vector< T, cms::cuda::HostAllocator< T >> &src, cudaStream_t stream)

bool	cudaCheck_ (const char file, int line, const char cmd, CUresult result, std::string_view description=std::string_view())

bool	cudaCheck_ (const char file, int line, const char cmd, cudaError_t result, std::string_view description=std::string_view())

int	currentDevice ()

allocator::GpuCachedBytes	deviceAllocatorStatus ()

int	deviceCount ()

bool	eventWorkHasCompleted (cudaEvent_t event)

	for (int i=first, nt=offsets[nh];i< nt;i+=gridDim.x *blockDim.x)

	for (auto pj=hist.begin(bs);pj< hist.end(be);++pj)

void	free_device (int device, void *ptr)

void	free_host (void *ptr)

EventCache &	getEventCache ()

StreamCache &	getStreamCache ()

void	launch (void(*kernel)(), LaunchParameters config)

template<typename F , typename... Args>
std::enable_if_t< std::is_void< std::result_of_t< F &&(Args &&...)> >::value >	launch (F *kernel, LaunchParameters config, Args &&... args)

void	launch_cooperative (void(*kernel)(), LaunchParameters config)

template<typename F , typename... Args>
std::enable_if_t< std::is_void< std::result_of_t< F &&(Args &&...)> >::value >	launch_cooperative (F *kernel, LaunchParameters config, Args &&... args)

template<typename T >
device::impl::make_device_unique_selector< T >::non_array	make_device_unique (cudaStream_t stream)

template<typename T >
device::impl::make_device_unique_selector< T >::unbounded_array	make_device_unique (size_t n, cudaStream_t stream)

template<typename T , typename... Args>
device::impl::make_device_unique_selector< T >::bounded_array	make_device_unique (Args &&...)=delete

template<typename T >
device::impl::make_device_unique_selector< T >::non_array	make_device_unique_uninitialized (cudaStream_t stream)

template<typename T >
device::impl::make_device_unique_selector< T >::unbounded_array	make_device_unique_uninitialized (size_t n, cudaStream_t stream)

template<typename T , typename... Args>
device::impl::make_device_unique_selector< T >::bounded_array	make_device_unique_uninitialized (Args &&...)=delete

template<typename T >
host::noncached::impl::make_host_unique_selector< T >::non_array	make_host_noncached_unique (unsigned int flags=cudaHostAllocDefault)

template<typename T >
host::noncached::impl::make_host_unique_selector< T >::unbounded_array	make_host_noncached_unique (size_t n, unsigned int flags=cudaHostAllocDefault)

template<typename T , typename... Args>
host::noncached::impl::make_host_unique_selector< T >::bounded_array	make_host_noncached_unique (Args &&...)=delete

template<typename T >
host::impl::make_host_unique_selector< T >::non_array	make_host_unique ()

template<typename T >
host::impl::make_host_unique_selector< T >::unbounded_array	make_host_unique (size_t n)

template<typename T >
host::impl::make_host_unique_selector< T >::non_array	make_host_unique (cudaStream_t stream)

template<typename T >
host::impl::make_host_unique_selector< T >::unbounded_array	make_host_unique (size_t n, cudaStream_t stream)

template<typename T , typename... Args>
host::impl::make_host_unique_selector< T >::bounded_array	make_host_unique (Args &&...)=delete

template<typename T >
host::impl::make_host_unique_selector< T >::non_array	make_host_unique_uninitialized (cudaStream_t stream)

template<typename T >
host::impl::make_host_unique_selector< T >::unbounded_array	make_host_unique_uninitialized (size_t n, cudaStream_t stream)

template<typename T , typename... Args>
host::impl::make_host_unique_selector< T >::bounded_array	make_host_unique_uninitialized (Args &&...)=delete

template<class T >
SimpleVector< T >	make_SimpleVector (int capacity, T *data)

template<class T >
SimpleVector< T > *	make_SimpleVector (SimpleVector< T > mem, int capacity, T data)

template<typename T >
void	memsetAsync (device::unique_ptr< T > &ptr, T value, cudaStream_t stream)

template<typename T >
void	memsetAsync (device::unique_ptr< T[]> &ptr, int value, size_t nelements, cudaStream_t stream)

int	numberOfDevices ()

bool	nvmlCheck_ (const char file, int line, const char cmd, nvmlReturn_t result, std::string_view description=std::string_view())

Variables
int	be = std::min(int(Hist::nbins() - 1), bs + n)

	bs = std::max(0, bs - n)

__host__ __device__ VT *	co

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V int Func	func

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int Histo::index_type *	mem

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V int	n

uint32_t	nh

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int	nthreads

uint32_t T const __restrict__ uint32_t const __restrict__	offsets

cudaStream_t int32_t	ONES

__host__ __device__ VT uint32_t	size

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int Histo::index_type cudaStream_t	stream

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t	totSize

uint32_t T const *__restrict__	v

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int Histo::index_type cudaStream_t	V

uint32_t T const __restrict__ uint32_t const __restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V	value

__host__ __device__ V V	wmax

__host__ __device__ V	wmin

Typedef Documentation

◆ SharedEventPtr

using cms::cuda::SharedEventPtr = typedef std::shared_ptr<std::remove_pointer_t<cudaEvent_t> >

Definition at line 14 of file SharedEventPtr.h.

◆ SharedStreamPtr

using cms::cuda::SharedStreamPtr = typedef std::shared_ptr<std::remove_pointer_t<cudaStream_t> >

Definition at line 14 of file SharedStreamPtr.h.

Function Documentation

◆ attribute() [1/4]

template<typename Histo , typename T >

cms::cuda::__attribute__ ( (always_inline) )

inline

◆ attribute() [2/4]

template<typename Assoc >

cms::cuda::__attribute__ ( (always_inline) )

inline

Definition at line 36 of file OneToManyAssoc.h.

References cms::cudacompat::__syncthreads(), assert(), cms::cudacompat::blockDim, cms::cudacompat::blockIdx, dqmdumpme::first, cms::cudacompat::gridDim, h, mps_fire::i, nt, and cms::cudacompat::threadIdx.

                                                                 {
       auto h = view.assoc;
       assert(1 == gridDim.x);
       assert(0 == blockIdx.x);
 
       int first = threadIdx.x;
 
       if (0 == first) {
         h->psws = 0;
         h->initStorage(view);
       }
       __syncthreads();
       for (int i = first, nt = h->totOnes(); i < nt; i += blockDim.x) {
         h->off[i] = 0;
       }
     }

◆ attribute() [3/4]

template<typename VT , typename T >

__host__ __device__ cms::cuda::__attribute__ ( (always_inline) ) const

inline

◆ attribute() [4/4]

__host__ __device__ cms::cuda::__attribute__ ( (always_inline) ) const &

inline

◆ abortOnCudaError()

void cms::cuda::abortOnCudaError	(	const char *	file,
		int	line,
		const char *	cmd,
		const char *	error,
		const char *	message,
		std::string_view	description = `std::string_view()`
	)

inline

Definition at line 21 of file cudaCheck.h.

References mps_setup::cmd, makeListRunsInFiles::description, relativeConstraints::error, geometryDiff::file, mps_splice::line, and MillePedeFileConverter_cfg::out.

Referenced by cudaCheck_().

                                                                                              {
       std::ostringstream out;
       out << "\n";
       out << file << ", line " << line << ":\n";
       out << "cudaCheck(" << cmd << ");\n";
       out << error << ": " << message << "\n";
       if (!description.empty())
         out << description << "\n";
       throw std::runtime_error(out.str());
     }

◆ abortOnNvmlError()

void cms::cuda::abortOnNvmlError	(	const char *	file,
		int	line,
		const char *	cmd,
		const char *	error,
		const char *	message,
		std::string_view	description = `std::string_view()`
	)

inline

Definition at line 20 of file nvmlCheck.h.

References mps_setup::cmd, makeListRunsInFiles::description, relativeConstraints::error, geometryDiff::file, mps_splice::line, and MillePedeFileConverter_cfg::out.

Referenced by nvmlCheck_().

                                                                                              {
       std::ostringstream out;
       out << "\n";
       out << file << ", line " << line << ":\n";
       out << "nvmlCheck(" << cmd << ");\n";
       out << error << ": " << message << "\n";
       if (!description.empty())
         out << description << "\n";
       throw std::runtime_error(out.str());
     }

◆ allocate_device()

void * cms::cuda::allocate_device	(	int	dev,
		size_t	nbytes,
		cudaStream_t	stream
	)

Definition at line 16 of file allocate_device.cc.

References cudaCheck, cms::cuda::allocator::getCachingDeviceAllocator(), stream, cond::impl::to_string(), UNLIKELY, and cms::cuda::allocator::useCaching.

Referenced by make_device_unique(), and make_device_unique_uninitialized().

                                                                      {
     void *ptr = nullptr;
     if constexpr (allocator::useCaching) {
       if (UNLIKELY(nbytes > maxAllocationSize)) {
         throw std::runtime_error("Tried to allocate " + std::to_string(nbytes) +
                                  " bytes, but the allocator maximum is " + std::to_string(maxAllocationSize));
       }
       cudaCheck(allocator::getCachingDeviceAllocator().DeviceAllocate(dev, &ptr, nbytes, stream));
     } else {
       ScopedSetDevice setDeviceForThisScope(dev);
       cudaCheck(cudaMalloc(&ptr, nbytes));
     }
     return ptr;
   }

◆ allocate_host()

void * cms::cuda::allocate_host	(	size_t	nbytes,
		cudaStream_t	stream
	)

Definition at line 15 of file allocate_host.cc.

References cudaCheck, cms::cuda::allocator::getCachingHostAllocator(), stream, cond::impl::to_string(), UNLIKELY, and cms::cuda::allocator::useCaching.

Referenced by make_host_unique(), and make_host_unique_uninitialized().

                                                           {
     void *ptr = nullptr;
     if constexpr (allocator::useCaching) {
       if (UNLIKELY(nbytes > maxAllocationSize)) {
         throw std::runtime_error("Tried to allocate " + std::to_string(nbytes) +
                                  " bytes, but the allocator maximum is " + std::to_string(maxAllocationSize));
       }
       cudaCheck(allocator::getCachingHostAllocator().HostAllocate(&ptr, nbytes, stream));
     } else {
       cudaCheck(cudaMallocHost(&ptr, nbytes));
     }
     return ptr;
   }

◆ assert()

cms::cuda::assert ( be >= bs )

◆ chooseDevice()

int cms::cuda::chooseDevice ( edm::StreamID id )

Definition at line 8 of file chooseDevice.cc.

References cms::Exception::addContext(), and or.

                                    {
     edm::Service<CUDAInterface> cuda;
     if (not cuda or not cuda->enabled()) {
       cms::Exception ex("CUDAError");
       ex << "Unable to choose current device because CUDAService is not preset or disabled.\n"
          << "If CUDAService was not explicitly disabled in the configuration, the probable\n"
          << "cause is that there is no GPU or there is some problem in the CUDA runtime or\n"
          << "drivers.";
       ex.addContext("Calling cms::cuda::chooseDevice()");
       throw ex;
     }
 
     // For startes we "statically" assign the device based on
     // edm::Stream number. This is suboptimal if the number of
     // edm::Streams is not a multiple of the number of CUDA devices
     // (and even then there is no load balancing).
     //
     // TODO: improve the "assignment" logic
     return id % cuda->numberOfDevices();
   }

◆ copyAsync() [1/6]

template<typename T >

void cms::cuda::copyAsync	(	device::unique_ptr< T > &	dst,
		const host::unique_ptr< T > &	src,
		cudaStream_t	stream
	)

inline

Definition at line 20 of file copyAsync.h.

References cudaCheck, math::cholesky::dst, TrackRefitter_38T_cff::src, stream, and relativeConstraints::value.

                                                                                                        {
       // Shouldn't compile for array types because of sizeof(T), but
       // let's add an assert with a more helpful message
       static_assert(std::is_array<T>::value == false,
                     "For array types, use the other overload with the size parameter");
       cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), sizeof(T), cudaMemcpyHostToDevice, stream));
     }

◆ copyAsync() [2/6]

template<typename T >

void cms::cuda::copyAsync	(	host::unique_ptr< T > &	dst,
		const device::unique_ptr< T > &	src,
		cudaStream_t	stream
	)

inline

Definition at line 38 of file copyAsync.h.

References cudaCheck, math::cholesky::dst, TrackRefitter_38T_cff::src, stream, and relativeConstraints::value.

                                                                                                        {
       static_assert(std::is_array<T>::value == false,
                     "For array types, use the other overload with the size parameter");
       cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), sizeof(T), cudaMemcpyDeviceToHost, stream));
     }

◆ copyAsync() [3/6]

template<typename T >

void cms::cuda::copyAsync	(	device::unique_ptr< T[]> &	dst,
		const host::unique_ptr< T[]> &	src,
		size_t	nelements,
		cudaStream_t	stream
	)

inline

Definition at line 47 of file copyAsync.h.

References cudaCheck, math::cholesky::dst, TrackRefitter_38T_cff::src, and stream.

                                                {
       cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), nelements * sizeof(T), cudaMemcpyHostToDevice, stream));
     }

◆ copyAsync() [4/6]

template<typename T >

void cms::cuda::copyAsync	(	host::unique_ptr< T[]> &	dst,
		const device::unique_ptr< T[]> &	src,
		size_t	nelements,
		cudaStream_t	stream
	)

inline

Definition at line 63 of file copyAsync.h.

References cudaCheck, math::cholesky::dst, TrackRefitter_38T_cff::src, and stream.

                                                {
       cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), nelements * sizeof(T), cudaMemcpyDeviceToHost, stream));
     }

◆ copyAsync() [5/6]

template<typename T >

void cms::cuda::copyAsync	(	cms::cuda::device::unique_ptr< T[]> &	dst,
		const std::vector< T, cms::cuda::HostAllocator< T >> &	src,
		cudaStream_t	stream
	)

inline

Definition at line 72 of file copyAsync.h.

References cudaCheck, math::cholesky::dst, TrackRefitter_38T_cff::src, and stream.

                                                {
       cudaCheck(cudaMemcpyAsync(dst.get(), src.data(), src.size() * sizeof(T), cudaMemcpyHostToDevice, stream));
     }

◆ copyAsync() [6/6]

template<typename T >

void cms::cuda::copyAsync	(	edm::propagate_const_array< cms::cuda::device::unique_ptr< T[]>> &	dst,
		const std::vector< T, cms::cuda::HostAllocator< T >> &	src,
		cudaStream_t	stream
	)

inline

Definition at line 80 of file copyAsync.h.

References cudaCheck, math::cholesky::dst, edm::get_underlying(), TrackRefitter_38T_cff::src, and stream.

                                                {
       cudaCheck(cudaMemcpyAsync(
           get_underlying(dst).get(), src.data(), src.size() * sizeof(T), cudaMemcpyHostToDevice, stream));
     }

◆ cudaCheck_() [1/2]

bool cms::cuda::cudaCheck_	(	const char *	file,
		int	line,
		const char *	cmd,
		CUresult	result,
		std::string_view	description = `std::string_view()`
	)

inline

Definition at line 37 of file cudaCheck.h.

References abortOnCudaError(), mps_setup::cmd, makeListRunsInFiles::description, relativeConstraints::error, geometryDiff::file, LIKELY, mps_splice::line, and mps_fire::result.

                                                                           {
       if (LIKELY(result == CUDA_SUCCESS))
         return true;
 
       const char* error = nullptr;
       const char* message = nullptr;
       cuGetErrorName(result, &error);
       cuGetErrorString(result, &message);
       abortOnCudaError(file, line, cmd, error, message, description);
       return false;
     }

◆ cudaCheck_() [2/2]

bool cms::cuda::cudaCheck_	(	const char *	file,
		int	line,
		const char *	cmd,
		cudaError_t	result,
		std::string_view	description = `std::string_view()`
	)

inline

Definition at line 53 of file cudaCheck.h.

References abortOnCudaError(), mps_setup::cmd, makeListRunsInFiles::description, relativeConstraints::error, geometryDiff::file, LIKELY, mps_splice::line, and mps_fire::result.

                                                                           {
       if (LIKELY(result == cudaSuccess))
         return true;
 
       const char* error = cudaGetErrorName(result);
       const char* message = cudaGetErrorString(result);
       abortOnCudaError(file, line, cmd, error, message, description);
       return false;
     }

◆ currentDevice()

int cms::cuda::currentDevice ( )

inline

Definition at line 10 of file currentDevice.h.

References cudaCheck.

Referenced by cms::cuda::ESProduct< HcalGainWidthsGPU::Product >::dataForCurrentDeviceAsync(), cms::cuda::StreamCache::get(), cms::cuda::EventCache::get(), make_device_unique(), make_device_unique_uninitialized(), and cms::cuda::allocator::minCachedBytes().

                                {
       int dev;
       cudaCheck(cudaGetDevice(&dev));
       return dev;
     }

◆ deviceAllocatorStatus()

allocator::GpuCachedBytes cms::cuda::deviceAllocatorStatus ( )

Definition at line 6 of file deviceAllocatorStatus.cc.

References notcub::CachingDeviceAllocator::CacheStatus(), and cms::cuda::allocator::getCachingDeviceAllocator().

6 { return allocator::getCachingDeviceAllocator().CacheStatus(); }

notcub::CachingDeviceAllocator::CacheStatus

GpuCachedBytes CacheStatus() const

Definition: CachingDeviceAllocator.h:728

cms::cuda::allocator::getCachingDeviceAllocator

notcub::CachingDeviceAllocator & getCachingDeviceAllocator()

Definition: getCachingDeviceAllocator.h:15

◆ deviceCount()

int cms::cuda::deviceCount ( )

inline

Definition at line 10 of file deviceCount.h.

References cudaCheck.

Referenced by cms::cuda::StreamCache::clear(), cms::cuda::EventCache::clear(), and cms::cuda::allocator::minCachedBytes().

                              {
       int ndevices;
       cudaCheck(cudaGetDeviceCount(&ndevices));
       return ndevices;
     }

◆ eventWorkHasCompleted()

bool cms::cuda::eventWorkHasCompleted ( cudaEvent_t event )

inline

Returns true if the work captured by the event (=queued to the CUDA stream at the point of cudaEventRecord()) has completed.

Returns false if any captured work is incomplete.

In case of errors, throws an exception.

Definition at line 18 of file eventWorkHasCompleted.h.

References cudaCheck, and runTheMatrix::ret.

Referenced by cms::cuda::ESProduct< HcalGainWidthsGPU::Product >::dataForCurrentDeviceAsync(), cms::cuda::EventCache::get(), and cms::cuda::ProductBase::isAvailable().

                                                          {
       const auto ret = cudaEventQuery(event);
       if (ret == cudaSuccess) {
         return true;
       } else if (ret == cudaErrorNotReady) {
         return false;
       }
       // leave error case handling to cudaCheck
       cudaCheck(ret);
       return false;  // to keep compiler happy
     }

◆ for() [1/2]

cms::cuda::for ( int i = first )

Definition at line 15 of file HistoContainer.h.

References assert(), mps_fire::i, nh, offsets, cuda_std::upper_bound(), and v.

Referenced by SiStripAPVRestorer::abnormalBaselineInspect(), l1tVertexFinder::AlgoSettings::AlgoSettings(), DumpFWRecoGeometry::analyze(), HGCalImagingAlgo::calculateDistanceToHigher(), spr::cGenSimInfo(), LHCInfoPerFill::cloneFill(), LHCInfo::cloneFill(), CandCombinerBase< OutputCollection, Cloner::CandPtr >::combine(), EERecHitGPU::convert_constant_data_(), HEFRecHitGPU::convert_constant_data_(), GsfBetheHeitlerUpdator::correctWeights(), CSCOfflineMonitor::doEfficiencies(), CSCOfflineMonitor::doStripDigis(), CSCOfflineMonitor::doWireDigis(), SiPixelStatusHarvester::dqmEndRun(), InvRingCalib::endOfLoop(), SiStripBadAPVAlgorithmFromClusterOccupancy::extractBadAPVs(), NtupleManager::FillFitParameters(), CSCEfficiency::fillStrips_info(), HLTCTPPSLocalTrackFilter::filter(), mkfit::MkBuilder::find_tracks_unroll_candidates(), ecaldqm::binning::getBinningEcal_(), HcalLutManager_test::getLutSetFromFile_test(), HcalDDDSimConstants::initialize(), FWGeometry::initMap(), InputTagDistributor::inputTag(), SiPixelTemplate::interpolate(), HGCalGeomParameters::loadGeometryHexagonModule(), SiStripFedCablingFakeESSource::make(), TSToSCAssociatorByEnergyScoreImpl::makeConnections(), MultiClusterAssociatorByEnergyScoreImpl::makeConnections(), LCToSCAssociatorByEnergyScoreImpl::makeConnections(), LCToCPAssociatorByEnergyScoreImpl::makeConnections(), TSToSimTSAssociatorByEnergyScoreImpl::makeConnections(), PSSDigitizerAlgorithm::module_killing_DB(), SSDigitizerAlgorithm::module_killing_DB(), operator<<(), operator>>(), SiPixelTemplateReco::PixelTempReco1D(), SiPixelTemplateReco2D::PixelTempReco2D(), SiPixelGenError::postInit(), SiPixelTemplate::postInit(), DAClusterizerInZ_vect::purge(), DAClusterizerInZT_vect::purge(), EcalPhiSymRecHitProducerBase::sumCache(), TPedValues::TPedValues(), EcalPedOffset::~EcalPedOffset(), SiStripLAProfileBooker::~SiStripLAProfileBooker(), and StraightTrackAlignment::~StraightTrackAlignment().

                                                                                  {
         auto off = cuda_std::upper_bound(offsets, offsets + nh + 1, i);
         assert((*off) > 0);
         int32_t ih = off - offsets - 1;
         assert(ih >= 0);
         assert(ih < int(nh));
         (*h).count(v[i], ih);
       }

◆ for() [2/2]

cms::cuda::for ( )

Definition at line 78 of file HistoContainer.h.

References func.

                                                               {
         func(*pj);
       }

◆ free_device()

void cms::cuda::free_device	(	int	device,
		void *	ptr
	)

Definition at line 31 of file allocate_device.cc.

References cudaCheck, cms::cuda::allocator::getCachingDeviceAllocator(), and cms::cuda::allocator::useCaching.

Referenced by cms::cuda::device::impl::DeviceDeleter::operator()().

                                           {
     if constexpr (allocator::useCaching) {
       cudaCheck(allocator::getCachingDeviceAllocator().DeviceFree(device, ptr));
     } else {
       ScopedSetDevice setDeviceForThisScope(device);
       cudaCheck(cudaFree(ptr));
     }
   }

◆ free_host()

void cms::cuda::free_host ( void * ptr )

Definition at line 29 of file allocate_host.cc.

References cudaCheck, cms::cuda::allocator::getCachingHostAllocator(), and cms::cuda::allocator::useCaching.

Referenced by cms::cuda::host::impl::HostDeleter::operator()().

                             {
     if constexpr (allocator::useCaching) {
       cudaCheck(allocator::getCachingHostAllocator().HostFree(ptr));
     } else {
       cudaCheck(cudaFreeHost(ptr));
     }
   }

◆ getEventCache()

EventCache & cms::cuda::getEventCache ( )

Definition at line 64 of file EventCache.cc.

References utilities::cache(), and CMS_THREAD_SAFE.

Referenced by CUDAService::CUDAService(), cms::cuda::ESProduct< HcalGainWidthsGPU::Product >::ESProduct(), and CUDAService::~CUDAService().

                               {
     // the public interface is thread safe
     CMS_THREAD_SAFE static EventCache cache;
     return cache;
   }

◆ getStreamCache()

StreamCache & cms::cuda::getStreamCache ( )

Definition at line 39 of file StreamCache.cc.

References utilities::cache(), and CMS_THREAD_SAFE.

Referenced by CUDAService::CUDAService(), TestCUDAAnalyzerGPU::endJob(), cms::cuda::impl::ScopedContextBase::ScopedContextBase(), TestCUDAAnalyzerGPU::TestCUDAAnalyzerGPU(), and CUDAService::~CUDAService().

                                 {
     // the public interface is thread safe
     CMS_THREAD_SAFE static StreamCache cache;
     return cache;
   }

◆ launch() [1/2]

void cms::cuda::launch	(	void(*)()	kernel,
		LaunchParameters	config
	)

inline

Definition at line 96 of file launch.h.

References cudaCheck.

                                                                   {
       cudaCheck(cudaLaunchKernel(
           (const void*)kernel, config.gridDim, config.blockDim, nullptr, config.sharedMem, config.stream));
     }

◆ launch() [2/2]

template<typename F , typename... Args>

std::enable_if_t<std::is_void<std::result_of_t<F && (Args && ...)> >::value> cms::cuda::launch	(	F *	kernel,
		LaunchParameters	config,
		Args &&...	args
	)

Definition at line 107 of file launch.h.

References writedatasetfile::args, cudaCheck, and size.

                                                                {
       using function_type = detail::kernel_traits<F>;
       typename function_type::argument_type_tuple args_copy(args...);
 
       constexpr auto size = function_type::arguments_size;
       void const* pointers[size];
 
       detail::pointer_setter<size>()(pointers, args_copy);
       cudaCheck(cudaLaunchKernel(
           (const void*)kernel, config.gridDim, config.blockDim, (void**)pointers, config.sharedMem, config.stream));
     }

◆ launch_cooperative() [1/2]

void cms::cuda::launch_cooperative	(	void(*)()	kernel,
		LaunchParameters	config
	)

inline

Definition at line 121 of file launch.h.

References cudaCheck.

                                                                               {
       cudaCheck(cudaLaunchCooperativeKernel(
           (const void*)kernel, config.gridDim, config.blockDim, nullptr, config.sharedMem, config.stream));
     }

◆ launch_cooperative() [2/2]

template<typename F , typename... Args>

std::enable_if_t<std::is_void<std::result_of_t<F && (Args && ...)> >::value> cms::cuda::launch_cooperative	(	F *	kernel,
		LaunchParameters	config,
		Args &&...	args
	)

Definition at line 132 of file launch.h.

References writedatasetfile::args, cudaCheck, and size.

                                                                            {
       using function_type = detail::kernel_traits<F>;
       typename function_type::argument_type_tuple args_copy(args...);
 
       constexpr auto size = function_type::arguments_size;
       void const* pointers[size];
 
       detail::pointer_setter<size>()(pointers, args_copy);
       cudaCheck(cudaLaunchCooperativeKernel(
           (const void*)kernel, config.gridDim, config.blockDim, (void**)pointers, config.sharedMem, config.stream));
     }

◆ make_device_unique() [1/3]

template<typename T >

device::impl::make_device_unique_selector<T>::non_array cms::cuda::make_device_unique ( cudaStream_t stream )

Definition at line 52 of file device_unique_ptr.h.

References allocate_device(), currentDevice(), mem, stream, and relativeConstraints::value.

Referenced by HBHERecHitProducerGPU::acquire(), HcalRawToDigiGPU::acquire(), and EcalRawToDigiGPU::acquire().

                                                                                                        {
       static_assert(std::is_trivially_constructible<T>::value,
                     "Allocating with non-trivial constructor on the device memory is not supported");
       int dev = currentDevice();
       void *mem = allocate_device(dev, sizeof(T), stream);
       return typename device::impl::make_device_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),
                                                                               device::impl::DeviceDeleter{dev}};
     }

◆ make_device_unique() [2/3]

template<typename T >

device::impl::make_device_unique_selector<T>::unbounded_array cms::cuda::make_device_unique	(	size_t	n,
		cudaStream_t	stream
	)

Definition at line 62 of file device_unique_ptr.h.

References allocate_device(), currentDevice(), mem, n, stream, and relativeConstraints::value.

                                                                                                                    {
       using element_type = typename std::remove_extent<T>::type;
       static_assert(std::is_trivially_constructible<element_type>::value,
                     "Allocating with non-trivial constructor on the device memory is not supported");
       int dev = currentDevice();
       void *mem = allocate_device(dev, n * sizeof(element_type), stream);
       return typename device::impl::make_device_unique_selector<T>::unbounded_array{
           reinterpret_cast<element_type *>(mem), device::impl::DeviceDeleter{dev}};
     }

◆ make_device_unique() [3/3]

template<typename T , typename... Args>

device::impl::make_device_unique_selector<T>::bounded_array cms::cuda::make_device_unique ( Args && ... )

delete

◆ make_device_unique_uninitialized() [1/3]

template<typename T >

device::impl::make_device_unique_selector<T>::non_array cms::cuda::make_device_unique_uninitialized ( cudaStream_t stream )

Definition at line 78 of file device_unique_ptr.h.

References allocate_device(), currentDevice(), mem, and stream.

                              {
       int dev = currentDevice();
       void *mem = allocate_device(dev, sizeof(T), stream);
       return typename device::impl::make_device_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),
                                                                               device::impl::DeviceDeleter{dev}};
     }

◆ make_device_unique_uninitialized() [2/3]

template<typename T >

device::impl::make_device_unique_selector<T>::unbounded_array cms::cuda::make_device_unique_uninitialized	(	size_t	n,
		cudaStream_t	stream
	)

Definition at line 87 of file device_unique_ptr.h.

References allocate_device(), currentDevice(), mem, n, and stream.

                                        {
       using element_type = typename std::remove_extent<T>::type;
       int dev = currentDevice();
       void *mem = allocate_device(dev, n * sizeof(element_type), stream);
       return typename device::impl::make_device_unique_selector<T>::unbounded_array{
           reinterpret_cast<element_type *>(mem), device::impl::DeviceDeleter{dev}};
     }

◆ make_device_unique_uninitialized() [3/3]

template<typename T , typename... Args>

device::impl::make_device_unique_selector<T>::bounded_array cms::cuda::make_device_unique_uninitialized ( Args && ... )

delete

◆ make_host_noncached_unique() [1/3]

template<typename T >

host::noncached::impl::make_host_unique_selector<T>::non_array cms::cuda::make_host_noncached_unique ( unsigned int flags = cudaHostAllocDefault )

The difference wrt. make_host_unique is that these do not cache, so they should not be called per-event.

Definition at line 47 of file host_noncached_unique_ptr.h.

References cudaCheck, HLT_2023v12_cff::flags, mem, and relativeConstraints::value.

                                                    {
       static_assert(std::is_trivially_constructible<T>::value,
                     "Allocating with non-trivial constructor on the pinned host memory is not supported");
       void *mem;
       cudaCheck(cudaHostAlloc(&mem, sizeof(T), flags));
       return typename host::noncached::impl::make_host_unique_selector<T>::non_array(reinterpret_cast<T *>(mem));
     }

◆ make_host_noncached_unique() [2/3]

template<typename T >

host::noncached::impl::make_host_unique_selector<T>::unbounded_array cms::cuda::make_host_noncached_unique	(	size_t	n,
		unsigned int	flags = `cudaHostAllocDefault`
	)

Definition at line 57 of file host_noncached_unique_ptr.h.

References cudaCheck, HLT_2023v12_cff::flags, mem, n, and relativeConstraints::value.

                                                              {
       using element_type = typename std::remove_extent<T>::type;
       static_assert(std::is_trivially_constructible<element_type>::value,
                     "Allocating with non-trivial constructor on the pinned host memory is not supported");
       void *mem;
       cudaCheck(cudaHostAlloc(&mem, n * sizeof(element_type), flags));
       return typename host::noncached::impl::make_host_unique_selector<T>::unbounded_array(
           reinterpret_cast<element_type *>(mem));
     }

◆ make_host_noncached_unique() [3/3]

template<typename T , typename... Args>

host::noncached::impl::make_host_unique_selector<T>::bounded_array cms::cuda::make_host_noncached_unique ( Args && ... )

delete

◆ make_host_unique() [1/5]

template<typename T >

host::impl::make_host_unique_selector<T>::non_array cms::cuda::make_host_unique ( )

Definition at line 63 of file host_unique_ptr.h.

References myMessageLogger_cff::alignment, cms::cuda::host::impl::kDefault, mem, size, and relativeConstraints::value.

Referenced by HcalRawToDigiGPU::acquire(), and EcalRawToDigiGPU::acquire().

                                                                               {
       static_assert(std::is_trivially_constructible<T>::value,
                     "Allocating with non-trivial constructor on the host memory is not supported");
       // Allocate a buffer aligned to 128 bytes, to match the CUDA cache line size
       const size_t alignment = 128;
       // std::aligned_alloc() requires the size to be a multiple of the alignment
       const size_t size = (sizeof(T) + alignment - 1) / alignment * alignment;
       void *mem = std::aligned_alloc(alignment, size);
       return typename host::impl::make_host_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),
                                                                           host::impl::MemoryType::kDefault};
     }

◆ make_host_unique() [2/5]

template<typename T >

host::impl::make_host_unique_selector<T>::unbounded_array cms::cuda::make_host_unique ( size_t n )

Definition at line 76 of file host_unique_ptr.h.

References myMessageLogger_cff::alignment, cms::cuda::host::impl::kDefault, mem, n, size, and relativeConstraints::value.

                                                                                             {
       using element_type = typename std::remove_extent<T>::type;
       static_assert(std::is_trivially_constructible<element_type>::value,
                     "Allocating with non-trivial constructor on the host memory is not supported");
       // Allocate a buffer aligned to 128 bytes, to match the CUDA cache line size
       const size_t alignment = 128;
       // std::aligned_alloc() requires the size to be a multiple of the alignment
       const size_t size = (n * sizeof(element_type) + alignment - 1) / alignment * alignment;
       void *mem = std::aligned_alloc(alignment, size);
       return typename host::impl::make_host_unique_selector<T>::unbounded_array{reinterpret_cast<element_type *>(mem),
                                                                                 host::impl::MemoryType::kDefault};
     }

◆ make_host_unique() [3/5]

template<typename T >

host::impl::make_host_unique_selector<T>::non_array cms::cuda::make_host_unique ( cudaStream_t stream )

Definition at line 91 of file host_unique_ptr.h.

References allocate_host(), cms::cuda::host::impl::kPinned, mem, stream, and relativeConstraints::value.

                                                                                                  {
       static_assert(std::is_trivially_constructible<T>::value,
                     "Allocating with non-trivial constructor on the host memory is not supported");
       void *mem = allocate_host(sizeof(T), stream);
       return typename host::impl::make_host_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),  //
                                                                           host::impl::MemoryType::kPinned};
     }

◆ make_host_unique() [4/5]

template<typename T >

host::impl::make_host_unique_selector<T>::unbounded_array cms::cuda::make_host_unique	(	size_t	n,
		cudaStream_t	stream
	)

Definition at line 100 of file host_unique_ptr.h.

References allocate_host(), cms::cuda::host::impl::kPinned, mem, n, stream, and relativeConstraints::value.

                                                                                                                  {
       using element_type = typename std::remove_extent<T>::type;
       static_assert(std::is_trivially_constructible<element_type>::value,
                     "Allocating with non-trivial constructor on the host memory is not supported");
       void *mem = allocate_host(n * sizeof(element_type), stream);
       return typename host::impl::make_host_unique_selector<T>::unbounded_array{reinterpret_cast<element_type *>(mem),
                                                                                 host::impl::MemoryType::kPinned};
     }

◆ make_host_unique() [5/5]

template<typename T , typename... Args>

host::impl::make_host_unique_selector<T>::bounded_array cms::cuda::make_host_unique ( Args && ... )

delete

◆ make_host_unique_uninitialized() [1/3]

template<typename T >

host::impl::make_host_unique_selector<T>::non_array cms::cuda::make_host_unique_uninitialized ( cudaStream_t stream )

Definition at line 115 of file host_unique_ptr.h.

References allocate_host(), cms::cuda::host::impl::kPinned, mem, and stream.

                                                                                                                {
       void *mem = allocate_host(sizeof(T), stream);
       return typename host::impl::make_host_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),  //
                                                                           host::impl::MemoryType::kPinned};
     }

◆ make_host_unique_uninitialized() [2/3]

template<typename T >

host::impl::make_host_unique_selector<T>::unbounded_array cms::cuda::make_host_unique_uninitialized	(	size_t	n,
		cudaStream_t	stream
	)

Definition at line 122 of file host_unique_ptr.h.

References allocate_host(), cms::cuda::host::impl::kPinned, mem, n, and stream.

                                        {
       using element_type = typename std::remove_extent<T>::type;
       void *mem = allocate_host(n * sizeof(element_type), stream);
       return typename host::impl::make_host_unique_selector<T>::unbounded_array{reinterpret_cast<element_type *>(mem),
                                                                                 host::impl::MemoryType::kPinned};
     }

◆ make_host_unique_uninitialized() [3/3]

template<typename T , typename... Args>

host::impl::make_host_unique_selector<T>::bounded_array cms::cuda::make_host_unique_uninitialized ( Args && ... )

delete

◆ make_SimpleVector() [1/2]

template<class T >

SimpleVector<T> cms::cuda::make_SimpleVector	(	int	capacity,
		T *	data
	)

Definition at line 124 of file SimpleVector.h.

References gpuVertexFinder::capacity(), data, and runTheMatrix::ret.

Referenced by SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA().

                                                              {
       SimpleVector<T> ret;
       ret.construct(capacity, data);
       return ret;
     }

◆ make_SimpleVector() [2/2]

template<class T >

SimpleVector<T>* cms::cuda::make_SimpleVector	(	SimpleVector< T > *	mem,
		int	capacity,
		T *	data
	)

Definition at line 132 of file SimpleVector.h.

References gpuVertexFinder::capacity(), data, mem, and runTheMatrix::ret.

                                                                                     {
       auto ret = new (mem) SimpleVector<T>();
       ret->construct(capacity, data);
       return ret;
     }

◆ memsetAsync() [1/2]

template<typename T >

void cms::cuda::memsetAsync	(	device::unique_ptr< T > &	ptr,
		T	value,
		cudaStream_t	stream
	)

inline

Definition at line 12 of file memsetAsync.h.

References cudaCheck, stream, relativeConstraints::value, and value.

Referenced by SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA().

                                                                                     {
       // Shouldn't compile for array types because of sizeof(T), but
       // let's add an assert with a more helpful message
       static_assert(std::is_array<T>::value == false,
                     "For array types, use the other overload with the size parameter");
       cudaCheck(cudaMemsetAsync(ptr.get(), value, sizeof(T), stream));
     }

◆ memsetAsync() [2/2]

template<typename T >

void cms::cuda::memsetAsync	(	device::unique_ptr< T[]> &	ptr,
		int	value,
		size_t	nelements,
		cudaStream_t	stream
	)

inline

The type of value is int because of cudaMemsetAsync() takes it as an int. Note that cudaMemsetAsync() sets the value of each byte to value. This may lead to unexpected results if sizeof(T) > 1 and value != 0.

Definition at line 27 of file memsetAsync.h.

References cudaCheck, stream, and value.

                                                                                                           {
       cudaCheck(cudaMemsetAsync(ptr.get(), value, nelements * sizeof(T), stream));
     }

◆ numberOfDevices()

int cms::cuda::numberOfDevices ( )

Definition at line 6 of file numberOfDevices.cc.

Referenced by cms::cuda::allocator::minCachedBytes().

                         {
     edm::Service<CUDAInterface> cuda;
     return (cuda and cuda->enabled()) ? cuda->numberOfDevices() : 0;
   }

◆ nvmlCheck_()

bool cms::cuda::nvmlCheck_	(	const char *	file,
		int	line,
		const char *	cmd,
		nvmlReturn_t	result,
		std::string_view	description = `std::string_view()`
	)

inline

Definition at line 36 of file nvmlCheck.h.

References abortOnNvmlError(), mps_setup::cmd, makeListRunsInFiles::description, relativeConstraints::error, geometryDiff::file, LIKELY, mps_splice::line, mps_fire::result, AlCaHLTBitMon_QueryRunRegistry::string, and cond::impl::to_string().

                                                                           {
       if (LIKELY(result == NVML_SUCCESS))
         return true;
 
       std::string error = "NVML Error " + std::to_string(result);
       const char* message = nvmlErrorString(result);
       abortOnNvmlError(file, line, cmd, error.c_str(), message, description);
       return false;
     }

Variable Documentation

◆ be

auto cms::cuda::be = std::min(int(Hist::nbins() - 1), bs + n)

Definition at line 75 of file HistoContainer.h.

Referenced by GEMGeometryBuilderFromCondDB::boundPlane(), ME0GeometryBuilderFromCondDB::build(), RPCGeometryBuilderFromCondDB::build(), GEMGeometryBuilderFromCondDB::buildEtaPartition(), GEMGeometryBuilder::buildEtaPartition(), RPCGeometryBuilder::buildGeometry(), FWLiteJobMetadataManager::doUpdate(), TFParams::inverpj(), FWEveViewManager::itemChanged(), BTagCalibrationReader::BTagCalibrationReaderImpl::load(), FWEveViewManager::modelChanges(), TrapezoidalPlaneBounds::TrapezoidalPlaneBounds(), gpuVertexFinder::while(), MuonsGrabber::writeDataForRelativeBX(), and LMFDat::writeDB().

◆ bs

cms::cuda::bs = std::max(0, bs - n)

Definition at line 76 of file HistoContainer.h.

◆ co

__host__ __device__ VT* cms::cuda::co

Definition at line 47 of file prefixScan.h.

Referenced by cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::__attribute__(), HLTExoticaPlotter::analyze(), HLTHiggsPlotter::analyze(), HLTHiggsSubAnalysis::analyze(), HLTExoticaSubAnalysis::analyze(), FWPFCandidateDetailView::build(), FlatHexagon::createCorners(), TruncatedPyramid::createCorners(), FlatTrd::createCorners(), BackToBackKinematicConstraint::derivative(), DDHtmlMaDetails::details(), perftools::EdmEventSize::dump(), perftools::detail::dump(), HGCalTBGeometry::get8Corners(), HGCalGeometry::get8Corners(), HGCalTBGeometry::getCorners(), HGCalGeometry::getCorners(), HGCalTBGeometry::getNewCorners(), HGCalGeometry::getNewCorners(), IdealObliquePrism::initCorners(), IdealZPrism::initCorners(), FlatHexagon::initCorners(), FlatTrd::initCorners(), CaloCellGeometry::inside(), makeEcalShape(), operator<<(), pat::PATVertexSlimmer::produce(), AlgorithmCalibration< T, CO >::readObject(), and BackToBackKinematicConstraint::value().

◆ func

__host__ __device__ V V Func const & cms::cuda::func

Initial value:

{

int bs = Hist::bin(value)

newFWLiteAna.bin

bin

Definition: newFWLiteAna.py:161

value

Definition: value.py:1

cms::cuda::bs

bs

Definition: HistoContainer.h:76

Definition at line 73 of file HistoContainer.h.

Referenced by for(), and cms::cuda::impl::ScopedContextHolderHelper::pushNextTask().

◆ mem

uint32_t T const* __restrict__ uint32_t const* __restrict__ int32_t int Histo::index_type* cms::cuda::mem

Definition at line 43 of file HistoContainer.h.

Referenced by make_device_unique(), make_device_unique_uninitialized(), make_host_noncached_unique(), make_host_unique(), make_host_unique_uninitialized(), and make_SimpleVector().

◆ n

uint32_t T const* __restrict__ uint32_t const* __restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V int cms::cuda::n

Definition at line 73 of file HistoContainer.h.

Referenced by cms::cuda::HostAllocator< T, FLAGS >::allocate(), make_device_unique(), make_device_unique_uninitialized(), make_host_noncached_unique(), make_host_unique(), and make_host_unique_uninitialized().

◆ nh

◆ nthreads

uint32_t T const* __restrict__ uint32_t const* __restrict__ int32_t int cms::cuda::nthreads

Definition at line 43 of file HistoContainer.h.

Referenced by addFiles().

◆ offsets

uint32_t T const *__restrict__ uint32_t const *__restrict__ cms::cuda::offsets

Initial value:

{

int first = blockDim.x * blockIdx.x + threadIdx.x

cms::cudacompat::threadIdx

const dim3 threadIdx

Definition: cudaCompat.h:29

cms::cudacompat::blockDim

const dim3 blockDim

Definition: cudaCompat.h:30

cms::cudacompat::blockIdx

const dim3 blockIdx

Definition: cudaCompat.h:32

dqmdumpme.first

first

Definition: dqmdumpme.py:55

Definition at line 13 of file HistoContainer.h.

Referenced by for().

◆ ONES

cudaStream_t int32_t cms::cuda::ONES

Definition at line 57 of file OneToManyAssoc.h.

Referenced by cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::ctNOnes().

◆ size

__host__ __device__ VT uint32_t cms::cuda::size

Definition at line 47 of file prefixScan.h.

Referenced by launch(), launch_cooperative(), and make_host_unique().

◆ stream

cudaStream_t cms::cuda::stream

Initial value:

= cudaStreamDefault
    ) {
      typename Histo::View view = {h, nullptr, mem, -1, totSize};
      launchZero(view, stream);
      countFromVector(h, nh, v, offsets);
      h->finalize();
      fillFromVector(h, nh, v, offsets);
    }
    
    template <typename Hist

Definition at line 51 of file HistoContainer.h.

◆ totSize

uint32_t T const* __restrict__ uint32_t const* __restrict__ int32_t cms::cuda::totSize

Definition at line 43 of file HistoContainer.h.

Referenced by hgcalsimclustertime::ComputeClusterTime::fixSizeHighestDensity().

◆ v

uint32_t T const *__restrict__ cms::cuda::v

Definition at line 11 of file HistoContainer.h.

Referenced by for(), cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::ilog2(), and cms::cuda::FlexiStorage< I, -1 >::init().

◆ V

uint32_t T const* __restrict__ uint32_t const* __restrict__ int32_t int Histo::index_type cudaStream_t cms::cuda::V

Definition at line 51 of file HistoContainer.h.

Referenced by TotemRPDQMSource::analyze(), L1TConfigDumper::analyze(), SiTrackerMultiRecHitUpdator::calcParameters(), Line::closerPointToLine(), Line::distance(), SignedImpactParameter3D::distanceWithJetAxis(), dot(), Chi2Strip1DEstimator::estimate(), MRHChi2MeasurementEstimator::estimate(), magfieldparam::HarmBasis3DCyl::EvalBphi(), magfieldparam::HarmBasis3DCyl::EvalRZ(), SurveyAlignmentPoints::findAlignPars(), CSCChamberFitter::fit(), hitfit::Chisq_Constrainer::fit(), TotemRPLocalTrackFitterAlgorithm::fitTrack(), SuepShower::generateFourVector(), MuonErrorMatrix::get(), MuonErrorMatrix::getFast(), ReferenceTrajectory::getHitProjectionMatrixT(), Rivet::HiggsTemplateCrossSections::getStage0Category(), Rivet::HiggsTemplateCrossSections::getStage1_1_Category(), Rivet::HiggsTemplateCrossSections::getStage1_1_Fine_Category(), Rivet::HiggsTemplateCrossSections::getStage1_2_Category(), Rivet::HiggsTemplateCrossSections::getStage1_2_Fine_Category(), Rivet::HiggsTemplateCrossSections::getStage1Category(), GlobalTrackerMuonAlignment::gradientGlobal(), GlobalTrackerMuonAlignment::gradientLocal(), PerigeeConversions::jacobianCurvilinear2Perigee(), PerigeeConversions::jacobianPerigee2Curvilinear(), IPTools::jetTrackDistance(), cmsutil::UnsortedDumbVectorMap< K, V >::operator[](), PFRecoTauTagInfoProducer::produce(), cond::payloadInspector::serializeValue(), TKinFitter::setCovMatrix(), IdealResult::solve(), edmtest::ViewAnalyzer::testProduct(), KFStrip1DUpdator::update(), and PosteriorWeightsCalculator::weights().

◆ value

uint32_t T const* __restrict__ uint32_t const* __restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V cms::cuda::value

Definition at line 73 of file HistoContainer.h.

Referenced by memsetAsync().

◆ wmax

__host__ __device__ V V cms::cuda::wmax

Definition at line 85 of file HistoContainer.h.

Referenced by mkfit::applyMaterialEffects(), HGCalDDDConstants::HGCalDDDConstants(), HGCalTBDDDConstants::HGCalTBDDDConstants(), PulseChiSqSNNLS::NNLS(), and MahiFit::nnls().

◆ wmin

__host__ __device__ V cms::cuda::wmin

Definition at line 85 of file HistoContainer.h.

Referenced by HGCalDDDConstants::HGCalDDDConstants(), and HGCalTBDDDConstants::HGCalTBDDDConstants().

Namespaces

Classes

Typedefs

Functions

Variables

Typedef Documentation

◆ SharedEventPtr

◆ SharedStreamPtr

Function Documentation

◆ __attribute__() [1/4]

◆ __attribute__() [2/4]

◆ __attribute__() [3/4]

◆ __attribute__() [4/4]

◆ abortOnCudaError()

◆ abortOnNvmlError()

◆ allocate_device()

◆ allocate_host()

◆ assert()

◆ chooseDevice()

◆ copyAsync() [1/6]

◆ copyAsync() [2/6]

◆ copyAsync() [3/6]

◆ copyAsync() [4/6]

◆ copyAsync() [5/6]

◆ copyAsync() [6/6]

◆ cudaCheck_() [1/2]

◆ cudaCheck_() [2/2]

◆ currentDevice()

◆ deviceAllocatorStatus()

◆ deviceCount()

◆ eventWorkHasCompleted()

◆ for() [1/2]

◆ for() [2/2]

◆ free_device()

◆ free_host()

◆ getEventCache()

◆ getStreamCache()

◆ launch() [1/2]

◆ launch() [2/2]

◆ launch_cooperative() [1/2]

◆ launch_cooperative() [2/2]

◆ make_device_unique() [1/3]

◆ make_device_unique() [2/3]

◆ make_device_unique() [3/3]

◆ make_device_unique_uninitialized() [1/3]

◆ make_device_unique_uninitialized() [2/3]

◆ make_device_unique_uninitialized() [3/3]

◆ make_host_noncached_unique() [1/3]

◆ make_host_noncached_unique() [2/3]

◆ make_host_noncached_unique() [3/3]

◆ make_host_unique() [1/5]

◆ make_host_unique() [2/5]

◆ make_host_unique() [3/5]

◆ make_host_unique() [4/5]

◆ make_host_unique() [5/5]

◆ make_host_unique_uninitialized() [1/3]

◆ make_host_unique_uninitialized() [2/3]

◆ make_host_unique_uninitialized() [3/3]

◆ make_SimpleVector() [1/2]

◆ make_SimpleVector() [2/2]

◆ memsetAsync() [1/2]

◆ memsetAsync() [2/2]

◆ numberOfDevices()

◆ nvmlCheck_()

Variable Documentation

◆ be

◆ bs

◆ co

◆ func

◆ mem

◆ n

◆ nh

◆ nthreads

◆ offsets

◆ ONES

◆ size

◆ stream

◆ totSize

◆ v

◆ V

◆ attribute() [1/4]

◆ attribute() [2/4]

◆ attribute() [3/4]

◆ attribute() [4/4]