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__(), atomicAdd_block(), cms::cuda::VecArray< TrackerTraits >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::extend(), gpuPixelDoublets::for(), gpuVertexFinder::for(), gpuClustering::for(), for(), caHitNtupletGeneratorKernels::for(), caHitNtupletGeneratorKernels::if(), cms::cuda::VecArray< TrackerTraits >::push_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::push_back(), and cms::cuda::SimpleVector< SiPixelErrorCompact >::shrink().

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

                              {
      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, SiStripPI::max, and runTheMatrix::ret.

Referenced by gpuPixelRecHits::__attribute__(), and atomicMax_block().

                              {
      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(), 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__(), atomicSub_block(), cms::cuda::VecArray< TrackerTraits >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::extend(), gpuPixelDoublets::for(), cms::cuda::VecArray< TrackerTraits >::push_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::push_back(), and cms::cuda::SimpleVector< SiPixelErrorCompact >::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(), and SimPFProducer::produce().

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(), and gpuCalibPixel::calibDigisPhase2().

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(), and gpuVertexFinder::while().