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

template<typename T >

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

inline

Definition at line 113 of file cudaCompat.h.

Referenced by TrackingRecHit2DSOAView::__attribute__(), SiPixelClustersCUDA::DeviceConstView::__attribute__(), SiPixelDigisCUDA::DeviceConstView::__attribute__(), calo::multifit::calculateChiSq(), gpuPixelDoublets::for(), pixelCPEforGPU::ParamsOnGPU::layer(), eigenSoA::ScalarSoA< uint8_t, S >::operator()(), eigenSoA::ScalarSoA< uint8_t, S >::operator[](), and calo::multifit::solve_forward_subst_matrix().

                               {
      return *x;
    }

gpuVertexFinder::__syncthreads ( )

inline

Definition at line 108 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuVertexFinder::__attribute__(), cms::cuda::__attribute__(), gpuClustering::for(), and for().

{}

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

inline

Definition at line 111 of file cudaCompat.h.

Referenced by gpuClustering::for().

{ return x; }

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

inline

Definition at line 110 of file cudaCompat.h.

{ return x; }

void cms::cudacompat::__threadfence ( )

inline

Definition at line 109 of file cudaCompat.h.

Referenced by GPUCACell::__attribute__().

{}

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< hindex_type, S+1, 5 *S >::__attribute__(), atomicAdd_block(), cms::cuda::VecArray< T, maxSize >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::extend(), gpuVertexFinder::for(), gpuClustering::for(), for(), gpuPixelDoublets::for(), if(), cms::cuda::VecArray< T, maxSize >::push_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::push_back(), and cms::cuda::SimpleVector< SiPixelErrorCompact >::shrink().

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

template<typename T1 , typename T2 >

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

Definition at line 68 of file cudaCompat.h.

References atomicAdd().

                                    {
      return atomicAdd(a, b);
    }

template<typename T1 , typename T2 >

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

Definition at line 36 of file cudaCompat.h.

Referenced by GPUCACell::__attribute__(), and atomicCAS_block().

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

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

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

template<typename T1 , typename T2 >

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

Definition at line 48 of file cudaCompat.h.

References a, and runTheMatrix::ret.

Referenced by gpuVertexFinder::__attribute__(), and atomicInc_block().

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

template<typename T1 , typename T2 >

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

Definition at line 56 of file cudaCompat.h.

References atomicInc().

                                    {
      return atomicInc(a, b);
    }

template<typename T1 , typename T2 >

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

Definition at line 97 of file cudaCompat.h.

References a, SiStripPI::max, and runTheMatrix::ret.

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

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

template<typename T1 , typename T2 >

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

Definition at line 104 of file cudaCompat.h.

References atomicMax().

                                    {
      return atomicMax(a, b);
    }

template<typename T1 , typename T2 >

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

Definition at line 85 of file cudaCompat.h.

References a, min(), and runTheMatrix::ret.

Referenced by gpuPixelRecHits::__attribute__(), atomicMin_block(), and gpuVertexFinder::while().

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

template<typename T1 , typename T2 >

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

Definition at line 92 of file cudaCompat.h.

References atomicMin().

                                    {
      return atomicMin(a, b);
    }

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< hindex_type, S+1, 5 *S >::__attribute__(), atomicSub_block(), cms::cuda::VecArray< T, maxSize >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::emplace_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::extend(), gpuPixelDoublets::for(), cms::cuda::VecArray< T, maxSize >::push_back(), cms::cuda::SimpleVector< SiPixelErrorCompact >::push_back(), and cms::cuda::SimpleVector< SiPixelErrorCompact >::shrink().

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

template<typename T1 , typename T2 >

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

Definition at line 80 of file cudaCompat.h.

References atomicSub().

                                    {
      return atomicSub(a, b);
    }

Variable Documentation

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

Definition at line 30 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuPixelDoublets::__attribute__(), gpuVertexFinder::__attribute__(), cms::cuda::__attribute__(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__(), gpuCalibPixel::calibDigis(), gpuClustering::for(), and gpuVertexFinder::while().

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

Definition at line 32 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuPixelDoublets::__attribute__(), cms::cuda::__attribute__(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__(), gpuCalibPixel::calibDigis(), gpuClustering::for(), and SimPFProducer::produce().

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

Definition at line 33 of file cudaCompat.h.

Referenced by gpuPixelDoublets::__attribute__(), cms::cuda::__attribute__(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__(), and gpuCalibPixel::calibDigis().

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

Definition at line 29 of file cudaCompat.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuPixelDoublets::__attribute__(), gpuVertexFinder::__attribute__(), cms::cuda::__attribute__(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__(), gpuCalibPixel::calibDigis(), gpuClustering::for(), gpuVertexFinder::if(), and gpuVertexFinder::while().