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

template<typename Hist , typename V , typename Func >
__host__ __device__	__attribute__ ((always_inline)) void forEachInWindow(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 (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

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

Definition at line 14 of file SharedEventPtr.h.

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

Definition at line 14 of file SharedStreamPtr.h.

Function Documentation

template<typename Histo , typename T >

uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ cms::cuda::__attribute__ ( (always_inline) )

inline

template<typename Assoc >

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

inline

Definition at line 36 of file OneToManyAssoc.h.

References cms::cudacompat::__syncthreads(), assert(), cms::cuda::OneToManyAssocView< Assoc >::assoc, cms::cudacompat::blockDim, cms::cudacompat::blockIdx, 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;
       }
     }

template<typename VT , typename T >

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

inline

template<typename Hist , typename V , typename Func >

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

inline

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 getPayloadData::description, and submitPVResolutionJobs::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());
     }

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 getPayloadData::description, and submitPVResolutionJobs::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());
     }

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

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

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

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

Definition at line 8 of file chooseDevice.cc.

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

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

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

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

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.

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

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.

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

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.

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

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, and edm::get_underlying().

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

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(), getPayloadData::description, relativeConstraints::error, and LIKELY.

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

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(), getPayloadData::description, and LIKELY.

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

int cms::cuda::currentDevice ( )

inline

Definition at line 10 of file currentDevice.h.

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

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

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

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

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
     }

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

Definition at line 15 of file HistoContainer.h.

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

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

cms::cuda::for ( )

Definition at line 78 of file HistoContainer.h.

References func.

                                                               {
         func(*pj);
       }

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

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

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

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

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

inline

Definition at line 96 of file launch.h.

References cms::cuda::LaunchParameters::blockDim, cudaCheck, cms::cuda::LaunchParameters::gridDim, cms::cuda::LaunchParameters::sharedMem, and cms::cuda::LaunchParameters::stream.

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

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, cms::cuda::LaunchParameters::blockDim, cudaCheck, cms::cuda::LaunchParameters::gridDim, cms::cuda::LaunchParameters::sharedMem, size, and cms::cuda::LaunchParameters::stream.

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

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

inline

Definition at line 121 of file launch.h.

References cms::cuda::LaunchParameters::blockDim, cudaCheck, cms::cuda::LaunchParameters::gridDim, cms::cuda::LaunchParameters::sharedMem, and cms::cuda::LaunchParameters::stream.

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

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, cms::cuda::LaunchParameters::blockDim, cudaCheck, cms::cuda::LaunchParameters::gridDim, cms::cuda::LaunchParameters::sharedMem, size, and cms::cuda::LaunchParameters::stream.

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

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, and relativeConstraints::value.

Referenced by HBHERecHitProducerGPU::acquire(), EcalRawToDigiGPU::acquire(), and HcalRawToDigiGPU::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}};
     }

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

template<typename T , typename... Args>

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

delete

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(), and mem.

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

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(), and mem.

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

template<typename T , typename... Args>

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

delete

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

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, mem, 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));
     }

template<typename T , typename... Args>

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

delete

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.

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

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

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

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.

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

template<typename T , typename... Args>

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

delete

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.

References allocate_host(), and mem.

Referenced by make_host_unique_uninitialized().

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

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.

References allocate_host(), make_host_unique_uninitialized(), and mem.

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

template<typename T , typename... Args>

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

delete

template<class T >

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

Definition at line 124 of file SimpleVector.h.

References cms::cuda::SimpleVector< T >::construct(), and runTheMatrix::ret.

Referenced by SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA().

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

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 mem, and runTheMatrix::ret.

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

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

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, and value.

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

int cms::cuda::numberOfDevices ( )

Definition at line 6 of file numberOfDevices.cc.

References fwrapper::cs, CUDAService::enabled(), and CUDAService::numberOfDevices().

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

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

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(), getPayloadData::description, LIKELY, 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

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

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

Definition at line 76 of file HistoContainer.h.

__host__ __device__ VT* cms::cuda::co

Definition at line 47 of file prefixScan.h.

Referenced by HLTExoticaPlotter::analyze(), HLTHiggsPlotter::analyze(), HLTHiggsSubAnalysis::analyze(), HLTExoticaSubAnalysis::analyze(), FWPFCandidateDetailView::build(), BackToBackKinematicConstraint::derivative(), DDHtmlMaDetails::details(), HGCalGeometry::get8Corners(), HGCalGeometry::getCorners(), HGCalGeometry::getNewCorners(), CaloCellGeometry::inside(), makeEcalShape(), pat::PATVertexSlimmer::produce(), AlgorithmCalibration< T, CO >::readObject(), and BackToBackKinematicConstraint::value().

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

Initial value:

{

int bs = Hist::bin(value)

relativeConstraints.value

tuple value

Definition: relativeConstraints.py:55

newFWLiteAna.bin

string bin

Definition: newFWLiteAna.py:161

cms::cuda::bs

bs

Definition: HistoContainer.h:76

Definition at line 73 of file HistoContainer.h.

Referenced by ntuplePrintersDiff::_mapdiff(), Config.Modifier::_toModify(), dqm::impl::MonitorElement::accessRootObject(), GeometryInterface::addExtractor(), CTPPSProtonReconstructionValidator::analyze(), FSQ::HandlerTemplate< TInputCandidateType, TOutputCandidateType, filter >::book(), BTagEntry::BTagEntry(), edm::stream::EDAnalyzerBase::callWhenNewProductsRegistered(), edm::EDAnalyzer::callWhenNewProductsRegistered(), edm::global::EDAnalyzerBase::callWhenNewProductsRegistered(), edm::limited::EDAnalyzerBase::callWhenNewProductsRegistered(), edm::one::EDAnalyzerBase::callWhenNewProductsRegistered(), edm::ProducerBase::callWhenNewProductsRegistered(), EcalTPGParamBuilder::computeLUT(), HFCherenkov::computeQEff(), BaseFunction::convertToArrays(), vertexPlots.VertexSummaryTable::create(), trackingPlots.TrackingSummaryTable::create(), GflashHadronShowerProfile::depthScale(), EcalZmassClient::dqmEndJob(), VVIObjFDetails::dzero(), VVIObjDetails::dzero(), sistripvvi::VVIObjDetails::dzero(), mkfit::ConfigJsonPatcher::exc_hdr(), PulseFitWithFunction::Fit_electronic(), SplitVertexResolution::fitResiduals(), TrackerOfflineValidationSummary::fitResiduals(), PVValHelper::fitResiduals(), TrackerOfflineValidation::fitResiduals(), SplitVertexResolution::fitResiduals_v0(), GflashHadronShowerProfile::fLnE1(), for(), GflashHadronShowerProfile::fTanh(), RPCSimAverage::getClSize(), RPCSimAverageNoise::getClSize(), RPCSimModelTiming::getClSize(), RPCSimAverageNoiseEff::getClSize(), RPCSimAverageNoiseEffCls::getClSize(), RPCSimAsymmetricCls::getClSize(), edm::refitem::GetRefPtrImpl< C, T, F, KEY >::getRefPtr_(), edm::refitem::GetRefPtrImpl< C, T, F, unsigned int >::getRefPtr_(), fwlite::DataGetterHelper::getThinnedKeyFrom(), BareRootProductGetter::getThinnedKeyFrom(), edm::EventPrincipal::getThinnedKeyFrom(), edm::ESProductHost< Product, RecordTypes >::ifRecordChanges(), edm::installSig(), main(), TFileDirectory::make(), edm::make_lambda_with_holder(), edm::make_waiting_task_with_holder(), edm::ESConsumesCollectorT< RECORD >::mayConsumeFrom(), operator<<(), JetResolution::parameterEtaEval(), EcalUncalibRecHitFixedAlphaBetaAlgo< C >::PerformAnalyticFit(), SymmetryFit::pol2_from_pol2(), SymmetryFit::pol2_from_pol3(), cms::cuda::impl::ScopedContextHolderHelper::pushNextTask(), stdcomb::recursive_combination(), edm::limited::impl::InputProcessBlockCacheHolder< T, CacheTypes >::registerProcessBlockCacheFiller(), edm::global::impl::InputProcessBlockCacheHolder< T, CacheTypes >::registerProcessBlockCacheFiller(), edm::one::impl::InputProcessBlockCacheHolder< T, CacheTypes >::registerProcessBlockCacheFiller(), cmsswSequenceInfo::serve(), cmsswConfigtrace::serve_main(), PSFitter::HybridMinimizer::SetFunction(), edm::ESConsumesCollectorT< RECORD >::setMayConsume(), reco::tau::sumPFVector(), PATStringObjectFunction::tryEval(), and PATStringCutObjectSelector::tryEval().

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

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

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

Definition at line 43 of file HistoContainer.h.

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

first

auto first

Definition: CAHitNtupletGeneratorKernelsImpl.h:127

cms::cudacompat::blockDim

const dim3 blockDim

Definition: cudaCompat.h:30

cms::cudacompat::blockIdx

const dim3 blockIdx

Definition: cudaCompat.h:32

Definition at line 13 of file HistoContainer.h.

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

__host__ __device__ VT uint32_t cms::cuda::size

Definition at line 47 of file prefixScan.h.

Referenced by launch(), and launch_cooperative().

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