cms::cudacompat Namespace Reference

Classes
struct	CPUTraits
struct	GPUTraits
struct	HostTraits

Functions
template<typename T >
T	__ldg (T const *x)
void	__syncthreads ()
bool	__syncthreads_and (bool x)
bool	__syncthreads_or (bool x)
void	__threadfence ()
template<typename T1 , typename T2 >
T1	atomicAdd (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicAdd_block (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicAnd (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicAnd_block (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicCAS (T1 *address, T1 compare, T2 val)
template<typename T1 , typename T2 >
T1	atomicCAS_block (T1 *address, T1 compare, T2 val)
template<typename T1 , typename T2 >
T1	atomicInc (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicInc_block (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicMax (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicMax_block (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicMin (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicMin_block (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicOr (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicOr_block (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicSub (T1 *a, T2 b)
template<typename T1 , typename T2 >
T1	atomicSub_block (T1 *a, T2 b)

Variables
const dim3	blockDim = {1, 1, 1}
const dim3	blockIdx = {0, 0, 0}
const dim3	gridDim = {1, 1, 1}
const dim3	threadIdx = {0, 0, 0}

Function Documentation

__ldg()

template<typename T >

T cms::cudacompat::__ldg ( T const *  x )

inline

Definition at line 137 of file cudaCompat.h.

References x.

Referenced by caHitNtupletGeneratorKernels::__attribute__(), calo::multifit::calculateChiSq(), gpuPixelDoublets::for(), pixelCPEforGPU::ParamsOnGPUT< TrackerTraits >::layer(), eigenSoA::ScalarSoA< T, S >::operator()(), eigenSoA::ScalarSoA< T, S >::operator[](), and calo::multifit::solve_forward_subst_matrix().

                               {
      return *x;
    }

__syncthreads()

gpuVertexFinder::__syncthreads ( )

inline

Definition at line 132 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), cms::cuda::__attribute__(), gpuClustering::for(), for(), and CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::merge_cands_for_seed().

{}

__syncthreads_and()

bool cms::cudacompat::__syncthreads_and ( bool  x )

inline

Definition at line 135 of file cudaCompat.h.

References x.

Referenced by gpuClustering::for().

{ return x; }

__syncthreads_or()

bool cms::cudacompat::__syncthreads_or ( bool  x )

inline

Definition at line 134 of file cudaCompat.h.

References x.

{ return x; }

__threadfence()

void cms::cudacompat::__threadfence ( )

inline

Definition at line 133 of file cudaCompat.h.

Referenced by GPUCACellT< TrackerTraits >::__attribute__().

{}

atomicAdd()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicAdd	(	T1 *	a,
		T2	b
	)

Definition at line 61 of file cudaCompat.h.

References a, b, and runTheMatrix::ret.

Referenced by gpuPixelRecHits::__attribute__(), cms::cuda::AtomicPairCounter::__attribute__(), cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::__attribute__(), cms::alpakatools::AtomicPairCounter::add(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::add(), atomicAdd_block(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::atomicIncrement(), cms::cuda::VecArray< T, maxSize >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::emplace_back(), cms::alpakatools::SimpleVector< T >::emplace_back(), cms::alpakatools::VecArray< TrackerTraits >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::extend(), cms::alpakatools::SimpleVector< T >::extend(), gpuPixelDoublets::for(), gpuVertexFinder::for(), gpuClustering::for(), for(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::for(), caHitNtupletGeneratorKernels::for(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::for(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_singleSeed(), caHitNtupletGeneratorKernels::if(), ALPAKA_ACCELERATOR_NAMESPACE::HGCalLayerClustersSoAAlgoKernelEnergy::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::raw::Kernel_unpack::operator()(), pixelClustering::ClusterChargeCut< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_BLFastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::LoadTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelRecHits::GetHits< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::HGCalLayerClustersSoAAlgoKernelPositionByHits::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_checkOverflows< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::HGCalLayerClustersSoAAlgoKernelPositionByHits2::operator()(), cms::alpakatools::multiBlockPrefixScan< T >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::FindClus< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::hcal::reconstruction::mahi::Kernel_prep1d_sameNumberOfSamples::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countSharedHit< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::SeedingTopoThresh::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TopoClusterContraction::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::FillRhfIndex::operator()(), cms::cuda::VecArray< T, maxSize >::push_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::push_back(), cms::alpakatools::SimpleVector< T >::push_back(), cms::alpakatools::VecArray< TrackerTraits >::push_back(), cms::alpakatools::radixSortImpl(), cms::cuda::SimpleVector< SiPixelErrorCompact >::shrink(), cms::alpakatools::SimpleVector< T >::shrink(), and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::sortByPt2().

                              {
      auto ret = *a;
      (*a) += b;
      return ret;
    }

atomicAdd_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicAdd_block	(	T1 *	a,
		T2	b
	)

Definition at line 68 of file cudaCompat.h.

References a, atomicAdd(), and b.

                                    {
      return atomicAdd(a, b);
    }

atomicAnd()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicAnd	(	T1 *	a,
		T2	b
	)

Definition at line 109 of file cudaCompat.h.

References a, b, and runTheMatrix::ret.

Referenced by atomicAnd_block().

                              {
      auto ret = *a;
      (*a) &= b;
      return ret;
    }

atomicAnd_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicAnd_block	(	T1 *	a,
		T2	b
	)

Definition at line 116 of file cudaCompat.h.

References a, atomicAnd(), and b.

                                    {
      return atomicAnd(a, b);
    }

atomicCAS()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicCAS	(	T1 *	address,
		T1	compare,
		T2	val
	)

Definition at line 36 of file cudaCompat.h.

References heppy_batch::val.

Referenced by GPUCACellT< TrackerTraits >::__attribute__(), atomicCAS_block(), atomicMaxF(), atomicMaxPair(), and gpuClustering::pixelStatus::promote().

                                                  {
      T1 old = *address;
      *address = old == compare ? val : old;
      return old;
    }

atomicCAS_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicCAS_block	(	T1 *	address,
		T1	compare,
		T2	val
	)

Definition at line 43 of file cudaCompat.h.

References atomicCAS(), and heppy_batch::val.

                                                        {
      return atomicCAS(address, compare, val);
    }

atomicInc()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicInc	(	T1 *	a,
		T2	b
	)

Definition at line 48 of file cudaCompat.h.

References a, b, and runTheMatrix::ret.

Referenced by atomicInc_block(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::clusterTracksByDensity(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::for(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitProducerKernelConstruct< CAL >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksDBSCAN::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksIterative::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::CountModules< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::FindClus< TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::RawToDigi_kernel< debug >::operator()().

                              {
      auto ret = *a;
      if ((*a) < T1(b))
        (*a)++;
      return ret;
    }

atomicInc_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicInc_block	(	T1 *	a,
		T2	b
	)

Definition at line 56 of file cudaCompat.h.

References a, atomicInc(), and b.

                                    {
      return atomicInc(a, b);
    }

atomicMax()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicMax	(	T1 *	a,
		T2	b
	)

Definition at line 97 of file cudaCompat.h.

References a, b, WZElectronSkims53X_cff::max, and runTheMatrix::ret.

Referenced by gpuPixelRecHits::__attribute__(), atomicMax_block(), ALPAKA_ACCELERATOR_NAMESPACE::pixelRecHits::GetHits< TrackerTraits >::operator()(), and cms::alpakatools::radixSortImpl().

                              {
      auto ret = *a;
      *a = std::max(*a, T1(b));
      return ret;
    }

atomicMax_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicMax_block	(	T1 *	a,
		T2	b
	)

Definition at line 104 of file cudaCompat.h.

References a, atomicMax(), and b.

                                    {
      return atomicMax(a, b);
    }

atomicMin()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicMin	(	T1 *	a,
		T2	b
	)

Definition at line 85 of file cudaCompat.h.

References a, b, SiStripPI::min, and runTheMatrix::ret.

Referenced by gpuPixelRecHits::__attribute__(), atomicMin_block(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::raw::Kernel_unpack::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksIterative::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelRecHits::GetHits< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::FindClus< TrackerTraits >::operator()(), and gpuVertexFinder::while().

                              {
      auto ret = *a;
      *a = std::min(*a, T1(b));
      return ret;
    }

atomicMin_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicMin_block	(	T1 *	a,
		T2	b
	)

Definition at line 92 of file cudaCompat.h.

References a, atomicMin(), and b.

                                    {
      return atomicMin(a, b);
    }

atomicOr()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicOr	(	T1 *	a,
		T2	b
	)

Definition at line 121 of file cudaCompat.h.

References a, b, and runTheMatrix::ret.

Referenced by atomicOr_block().

                             {
      auto ret = *a;
      (*a) |= b;
      return ret;
    }

atomicOr_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicOr_block	(	T1 *	a,
		T2	b
	)

Definition at line 128 of file cudaCompat.h.

References a, atomicOr(), and b.

                                   {
      return atomicOr(a, b);
    }

atomicSub()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicSub	(	T1 *	a,
		T2	b
	)

Definition at line 73 of file cudaCompat.h.

References a, b, and runTheMatrix::ret.

Referenced by cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::__attribute__(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::atomicDecrement(), atomicSub_block(), cms::cuda::VecArray< T, maxSize >::emplace_back(), cms::alpakatools::SimpleVector< T >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::emplace_back(), cms::alpakatools::VecArray< TrackerTraits >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::extend(), cms::alpakatools::SimpleVector< T >::extend(), gpuPixelDoublets::for(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::for(), cms::cuda::VecArray< T, maxSize >::push_back(), cms::alpakatools::SimpleVector< T >::push_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::push_back(), cms::alpakatools::VecArray< TrackerTraits >::push_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::shrink(), and cms::alpakatools::SimpleVector< T >::shrink().

                              {
      auto ret = *a;
      (*a) -= b;
      return ret;
    }

atomicSub_block()

template<typename T1 , typename T2 >

T1 cms::cudacompat::atomicSub_block	(	T1 *	a,
		T2	b
	)

Definition at line 80 of file cudaCompat.h.

References a, atomicSub(), and b.

                                    {
      return atomicSub(a, b);
    }

Variable Documentation

blockDim

const dim3 cms::cudacompat::blockDim = {1, 1, 1}

Definition at line 30 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), cms::cuda::__attribute__(), caHitNtupletGeneratorKernels::__attribute__(), cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::__attribute__(), cms::rocmtest::add_vectors_d(), cms::rocmtest::add_vectors_f(), gpuCalibPixel::calibDigis(), gpuCalibPixel::calibDigisPhase2(), gpuClustering::for(), gpuPixelDoublets::for(), and gpuVertexFinder::while().

blockIdx

const dim3 cms::cudacompat::blockIdx = {0, 0, 0}

Definition at line 32 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), cms::cuda::__attribute__(), caHitNtupletGeneratorKernels::__attribute__(), cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::__attribute__(), cms::rocmtest::add_vectors_d(), cms::rocmtest::add_vectors_f(), gpuCalibPixel::calibDigis(), gpuCalibPixel::calibDigisPhase2(), gpuClustering::for(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_exotic(), cms::alpakatools::multiBlockPrefixScan< T >::operator()(), SimPFProducer::produce(), and cms::alpakatools::radixSortMulti().

gridDim

const dim3 cms::cudacompat::gridDim = {1, 1, 1}

Definition at line 33 of file cudaCompat.h.

Referenced by cms::cuda::__attribute__(), caHitNtupletGeneratorKernels::__attribute__(), cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::__attribute__(), cms::rocmtest::add_vectors_d(), cms::rocmtest::add_vectors_f(), gpuCalibPixel::calibDigis(), gpuCalibPixel::calibDigisPhase2(), and ALPAKA_ACCELERATOR_NAMESPACE::ecal::raw::Kernel_unpack::operator()().

threadIdx

const dim3 cms::cudacompat::threadIdx = {0, 0, 0}

Definition at line 29 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), cms::cuda::__attribute__(), caHitNtupletGeneratorKernels::__attribute__(), cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >::__attribute__(), cms::rocmtest::add_vectors_d(), cms::rocmtest::add_vectors_f(), gpuCalibPixel::calibDigis(), gpuCalibPixel::calibDigisPhase2(), gpuClustering::for(), gpuVertexFinder::if(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::raw::Kernel_unpack::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_BLFastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_FastFit< N, TrackerTraits >::operator()(), cms::alpakatools::multiBlockPrefixScan< T >::operator()(), and gpuVertexFinder::while().