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	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 VT , typename T >
__host__ __device__	__attribute__ ((always_inline)) void blockPrefixScan(VT const *ci

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

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

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

void	abortOnCudaError (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 (cms::cuda::device::unique_ptr< T[]> &dst, const std::vector< T, cms::cuda::HostAllocator< T >> &src, cudaStream_t stream)

template<typename T >
void	copyAsync (device::unique_ptr< T > &dst, const host::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 (edm::propagate_const_array< cms::cuda::device::unique_ptr< T[]>> &dst, const std::vector< T, cms::cuda::HostAllocator< 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 (host::unique_ptr< T[]> &dst, const device::unique_ptr< T[]> &src, size_t nelements, cudaStream_t stream)

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

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

int	currentDevice ()

allocator::GpuCachedBytes	deviceAllocatorStatus ()

int	deviceCount ()

bool	eventWorkHasCompleted (cudaEvent_t event)

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

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

void	free_device (int device, void *ptr)

void	free_host (void *ptr)

EventCache &	getEventCache ()

StreamCache &	getStreamCache ()

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 (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)

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

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 (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_uninitialized (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>
host::noncached::impl::make_host_unique_selector< T >::bounded_array	make_host_noncached_unique (Args &&...)=delete

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 >
host::noncached::impl::make_host_unique_selector< T >::non_array	make_host_noncached_unique (unsigned int flags=cudaHostAllocDefault)

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 (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_uninitialized (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<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 ()

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

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

__host__ __device__ VT *	co

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

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

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

uint32_t	nh

uint32_t const T __restrict__ const uint32_t __restrict__ int32_t int	nthreads

uint32_t const T __restrict__ const uint32_t __restrict__	offsets

cudaStream_t int32_t	ONES

__host__ __device__ VT uint32_t	size

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

uint32_t const T __restrict__ const uint32_t __restrict__ int32_t	totSize

uint32_t const T *__restrict__	v

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

uint32_t const T __restrict__ const uint32_t __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 VT , typename T >

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

inline

◆ attribute() [2/4]

template<typename Histo , typename T >

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

inline

◆ attribute() [3/4]

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

inline

◆ attribute() [4/4]

template<typename Assoc >

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

inline

Definition at line 36 of file OneToManyAssoc.h.

                                                                 {
       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;
       }
     }

References cms::cudacompat::__syncthreads(), assert(), cms::cuda::OneToManyAssocView< Assoc >::assoc, cms::cudacompat::blockDim, cms::cudacompat::blockIdx, first, cms::cudacompat::gridDim, mps_fire::i, nt, and cms::cudacompat::threadIdx.

◆ 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.

                                                                                              {
       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());
     }

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

Referenced by cudaCheck_().

◆ allocate_device()

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

Definition at line 16 of file allocate_device.cc.

                                                                      {
     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;
   }

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

Referenced by make_device_unique(), and make_device_unique_uninitialized().

◆ allocate_host()

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

Definition at line 15 of file allocate_host.cc.

                                                           {
     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;
   }

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

Referenced by make_host_unique(), and make_host_unique_uninitialized().

◆ assert()

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

◆ chooseDevice()

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

Definition at line 8 of file chooseDevice.cc.

                                    {
     edm::Service<CUDAService> cudaService;
     if (not cudaService->enabled()) {
       cms::Exception ex("CUDAError");
       ex << "Unable to choose current device because CUDAService is disabled. If CUDAService was not explicitly\n"
             "disabled in the configuration, the probable cause is that there is no GPU or there is some problem\n"
             "in the CUDA runtime or 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 % cudaService->numberOfDevices();
   }

References cms::Exception::addContext(), CUDAService::enabled(), and CUDAService::numberOfDevices().

◆ copyAsync() [1/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.

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

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

◆ copyAsync() [2/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.

                                                                                                        {
       // 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));
     }

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

◆ 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.

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

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

◆ copyAsync() [4/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.

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

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

◆ copyAsync() [5/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.

                                                                                                        {
       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));
     }

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

◆ copyAsync() [6/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.

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

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

◆ cudaCheck_() [1/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.

                                                                           {
       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;
     }

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

◆ cudaCheck_() [2/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.

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

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

◆ currentDevice()

int cms::cuda::currentDevice ( )

inline

Definition at line 10 of file currentDevice.h.

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

References cudaCheck.

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

◆ deviceAllocatorStatus()

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

Definition at line 6 of file deviceAllocatorStatus.cc.

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

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

◆ deviceCount()

int cms::cuda::deviceCount ( )

inline

Definition at line 10 of file deviceCount.h.

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

References cudaCheck.

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

◆ 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.

                                                          {
       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
     }

References cudaCheck, and runTheMatrix::ret.

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

◆ for() [1/2]

cms::cuda::for ( )

Definition at line 78 of file HistoContainer.h.

                                                               {
         func(*pj);
       }

References func.

◆ for() [2/2]

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

Definition at line 15 of file HistoContainer.h.

                                                                                  {
         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);
       }

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(), LHCInfo::cloneFill(), CandCombinerBase< OutputCollection, Cloner::CandPtr >::combine(), EERecHitGPU::convert_constant_data_(), HEFRecHitGPU::convert_constant_data_(), GsfBetheHeitlerUpdator::correctWeights(), CSCOfflineMonitor::doEfficiencies(), CSCValidation::doEfficiencies(), CSCValidation::doGasGain(), CSCValidation::doPedestalNoise(), CSCOfflineMonitor::doStripDigis(), CSCValidation::doStripDigis(), CSCOfflineMonitor::doWireDigis(), CSCValidation::doWireDigis(), SiPixelStatusHarvester::dqmEndRun(), InvRingCalib::endOfLoop(), SiStripBadAPVAlgorithmFromClusterOccupancy::extractBadAPVs(), NtupleManager::FillFitParameters(), CSCEfficiency::fillStrips_info(), HLTCTPPSLocalTrackFilter::filter(), ecaldqm::binning::getBinningEcal_(), HcalLutManager_test::getLutSetFromFile_test(), HcalDDDSimConstants::initialize(), FWGeometry::initMap(), InputTagDistributor::inputTag(), SiPixelTemplate::interpolate(), HGCalGeomParameters::loadGeometryHexagonModule(), SiStripFedCablingFakeESSource::make(), TSToSCAssociatorByEnergyScoreImpl::makeConnections(), TSToSimTSAssociatorByEnergyScoreImpl::makeConnections(), LCToCPAssociatorByEnergyScoreImpl::makeConnections(), MultiClusterAssociatorByEnergyScoreImpl::makeConnections(), LCToSCAssociatorByEnergyScoreImpl::makeConnections(), HcalFlexiHardcodeGeometryLoader::makeHECells(), operator<<(), operator>>(), SiPixelTemplateReco::PixelTempReco1D(), SiPixelTemplateReco2D::PixelTempReco2D(), SiPixelGenError::postInit(), SiPixelTemplate::postInit(), DAClusterizerInZ_vect::purge(), DAClusterizerInZT_vect::purge(), TPedValues::TPedValues(), EcalPedOffset::~EcalPedOffset(), SiStripLAProfileBooker::~SiStripLAProfileBooker(), and StraightTrackAlignment::~StraightTrackAlignment().

◆ free_device()

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

Definition at line 31 of file allocate_device.cc.

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

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

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

◆ free_host()

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

Definition at line 29 of file allocate_host.cc.

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

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

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

◆ getEventCache()

EventCache & cms::cuda::getEventCache ( )

Definition at line 64 of file EventCache.cc.

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

References utilities::cache(), and CMS_THREAD_SAFE.

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

◆ getStreamCache()

StreamCache & cms::cuda::getStreamCache ( )

Definition at line 39 of file StreamCache.cc.

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

References utilities::cache(), and CMS_THREAD_SAFE.

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

◆ launch() [1/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.

                                                                {
       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));
     }

References writedatasetfile::args, cudaCheck, and size.

◆ launch() [2/2]

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

inline

Definition at line 96 of file launch.h.

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

References cudaCheck.

◆ launch_cooperative() [1/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.

                                                                            {
       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));
     }

References writedatasetfile::args, cudaCheck, and size.

◆ launch_cooperative() [2/2]

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

inline

Definition at line 121 of file launch.h.

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

References cudaCheck.

◆ make_device_unique() [1/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() [2/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.

                                                                                                        {
       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}};
     }

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

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

◆ make_device_unique() [3/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.

                                                                                                                    {
       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}};
     }

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

◆ make_device_unique_uninitialized() [1/3]

template<typename T , typename... Args>

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

delete

◆ make_device_unique_uninitialized() [2/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.

                              {
       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}};
     }

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

◆ make_device_unique_uninitialized() [3/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.

                                        {
       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}};
     }

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

◆ make_host_noncached_unique() [1/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_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.

                                                              {
       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));
     }

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

◆ make_host_noncached_unique() [3/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.

                                                    {
       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));
     }

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

◆ make_host_unique() [1/3]

template<typename T , typename... Args>

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

delete

◆ make_host_unique() [2/3]

template<typename T >

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

Definition at line 41 of file host_unique_ptr.h.

                                                                                                  {
       static_assert(std::is_trivially_constructible<T>::value,
                     "Allocating with non-trivial constructor on the pinned 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)};
     }

References allocate_host(), mem, stream, and relativeConstraints::value.

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

◆ make_host_unique() [3/3]

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 49 of file host_unique_ptr.h.

                                                                                                                  {
       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 = allocate_host(n * sizeof(element_type), stream);
       return typename host::impl::make_host_unique_selector<T>::unbounded_array{reinterpret_cast<element_type *>(mem)};
     }

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

◆ make_host_unique_uninitialized() [1/3]

template<typename T , typename... Args>

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

delete

◆ make_host_unique_uninitialized() [2/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 62 of file host_unique_ptr.h.

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

References allocate_host(), mem, and stream.

◆ make_host_unique_uninitialized() [3/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 68 of file host_unique_ptr.h.

                                        {
       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)};
     }

References allocate_host(), mem, n, and stream.

◆ 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.

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

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

Referenced by SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA().

◆ 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.

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

References gpuVertexFinder::capacity(), data, mem, and runTheMatrix::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.

                                                                                     {
       // 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));
     }

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

Referenced by SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA().

◆ 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.

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

References cudaCheck, stream, and value.

◆ numberOfDevices()

int cms::cuda::numberOfDevices ( )

Definition at line 6 of file numberOfDevices.cc.

                         {
     edm::Service<CUDAService> cs;
     return cs->enabled() ? cs->numberOfDevices() : 0;
   }

References fwrapper::cs.

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

Variable Documentation

◆ be

◆ 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< hindex_type, S+1, 5 *S >::__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(), HGCalGeometry::get8Corners(), HGCalGeometry::getCorners(), 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 const Func & cms::cuda::func

Initial value:

{

int bs = Hist::bin(value)

Definition at line 73 of file HistoContainer.h.

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

◆ mem

uint32_t const T* __restrict__ const uint32_t* __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 const T* __restrict__ const uint32_t* __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

uint32_t cms::cuda::nh

Definition at line 11 of file HistoContainer.h.

Referenced by CTPPSPixelDQMSource::analyze(), HcalTB06Analysis::analyze(), HcalHitValidation::analyzeLayer(), gen::HydjetHadronizer::call_hyinit(), for(), cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::histOff(), TrackListCombiner::produce(), EcalTBHodoscopeRecInfoAlgo::reconstruct(), PHcalTB06Info::saveHit(), PHcalTB04Info::saveHit(), and jsoncollector::DataPoint::snap().

◆ nthreads

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

Definition at line 43 of file HistoContainer.h.

Referenced by addFiles().

◆ offsets

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

Initial value:

{

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

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< hindex_type, S+1, 5 *S >::ctNOnes().

◆ size

__host__ __device__ VT uint32_t cms::cuda::size

Definition at line 47 of file prefixScan.h.

Referenced by launch(), and launch_cooperative().

◆ 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 const T* __restrict__ const uint32_t* __restrict__ int32_t cms::cuda::totSize

Definition at line 43 of file HistoContainer.h.

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

◆ v

uint32_t const T *__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 const T* __restrict__ const uint32_t* __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(), tauImpactParameter::ParticleBuilder::createLorentzVectorParticle(), 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(), MuonErrorMatrix::get(), MuonErrorMatrix::getFast(), ReferenceTrajectory::getHitProjectionMatrixT(), GlobalTrackerMuonAlignment::gradientGlobal(), GlobalTrackerMuonAlignment::gradientLocal(), PerigeeConversions::jacobianCurvilinear2Perigee(), PerigeeConversions::jacobianPerigee2Curvilinear(), IPTools::jetTrackDistance(), cmsutil::UnsortedDumbVectorMap< K, V >::operator[](), PFRecoTauTagInfoProducer::produce(), cond::payloadInspector::serializeValue(), TKinFitter::setCovMatrix(), IdealResult::solve(), KFStrip1DUpdator::update(), and PosteriorWeightsCalculator::weights().

◆ value

uint32_t const T* __restrict__ const uint32_t* __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 HGCalDDDConstants::HGCalDDDConstants(), PulseChiSqSNNLS::NNLS(), and MahiFit::nnls().

◆ wmin

__host__ __device__ V cms::cuda::wmin

Definition at line 85 of file HistoContainer.h.

Referenced by HGCalDDDConstants::HGCalDDDConstants().

Namespaces

Classes

Typedefs

Functions

Variables

Typedef Documentation

◆ SharedEventPtr

◆ SharedStreamPtr

Function Documentation

◆ __attribute__() [1/4]

◆ __attribute__() [2/4]

◆ __attribute__() [3/4]

◆ __attribute__() [4/4]

◆ abortOnCudaError()

◆ 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/3]

◆ make_host_unique() [2/3]

◆ make_host_unique() [3/3]

◆ 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()

Variable Documentation

◆ be

◆ bs

◆ co

◆ func

◆ mem

◆ n

◆ nh

◆ nthreads

◆ offsets

◆ ONES

◆ size

◆ stream

◆ totSize

◆ v

◆ V

◆ value

◆ wmax

◆ wmin

◆ attribute() [1/4]

◆ attribute() [2/4]

◆ attribute() [3/4]

◆ attribute() [4/4]