gpuMCMCProcessorUtils.cu File Reference

#include "Fitters/gpuMCMCProcessorUtils.cuh"

Include dependency graph for gpuMCMCProcessorUtils.cu:

Go to the source code of this file.

Functions
__global__ void	EvalOnGPU_AutoCorr (const float *__restrict__ ParStep_gpu, const float *__restrict__ ParamSums_gpu, float *NumeratorSum_gpu, float *DenomSum_gpu)
	Eval autocorrelations based on Box and Jenkins. More...

Variables
__device__ __constant__ int	d_nLag
__device__ __constant__ int	d_nDraws
__device__ __constant__ int	d_nEntries

Function Documentation

EvalOnGPU_AutoCorr()

__global__ void EvalOnGPU_AutoCorr	(	const float *__restrict__	ParStep_gpu,
		const float *__restrict__	ParamSums_gpu,
		float *	NumeratorSum_gpu,
		float *	DenomSum_gpu
	)

Eval autocorrelations based on Box and Jenkins.

Parameters

ParStep_gpu	Parameter value at each step
NumeratorSum_gpu	Sum used for nominator of autocorrelation calculations
ParamSums_gpu	Overall sum for each parameter over all steps
DenomSum_gpu	Sum used for denominator of autocorrelation calculations

Definition at line 109 of file gpuMCMCProcessorUtils.cu.

                          {
//*********************************************************
  const unsigned int CurrentLagNum = (blockIdx.x * blockDim.x + threadIdx.x);
 
  //KS: Accessing shared memory is much much faster than global memory hence we use shared memory for calculation and then write to global memory
  __shared__ float shared_NumeratorSum[_BlockSize_];
  __shared__ float shared_DenomSum[_BlockSize_];
 
  // this is the stopping condition!
  if (CurrentLagNum < d_nLag*d_nDraws)
  {
      shared_NumeratorSum[threadIdx.x] = 0;
      shared_DenomSum[threadIdx.x] = 0;
 
      //KS: Might consider caching this information, which MIGHT be faster too lazy right now
      const int Param = int(CurrentLagNum/d_nLag);
      const int nLag = CurrentLagNum - Param*d_nLag;
      // Loop over the number of entries
      for (int i = 0; i < d_nEntries; ++i)
      {
        //KS: Use fmaf to have it tiny bit faster, for something easier to read: Param*d_nEntries + i
        int CurrParStep = fmaf(Param, d_nEntries, i);
        const float Diff = ParStep_gpu[CurrParStep]-ParamSums_gpu[Param];
        // Only sum the numerator up to i = N-k
        if (i < d_nEntries-nLag)
        {
          //KS: Use fmaf to have it tiny bit faster, for something easier to read: Param*d_nEntries + (i + nLag)
          CurrParStep = fmaf(Param, d_nEntries, i + nLag);
          const float LagTerm = ParStep_gpu[CurrParStep]-ParamSums_gpu[Param];
          const float Product = Diff*LagTerm;
          shared_NumeratorSum[threadIdx.x] += Product;
        }
        // Square the difference to form the denominator
        const float Denom = Diff*Diff;
        shared_DenomSum[threadIdx.x] += Denom;
      }
 
      //KS: Make sure threads are synchronised before moving to global memory
      __syncthreads();
      NumeratorSum_gpu[CurrentLagNum] = shared_NumeratorSum[threadIdx.x];
      DenomSum_gpu[CurrentLagNum]     = shared_DenomSum[threadIdx.x];
  }
}

Variable Documentation

d_nDraws

__device__ __constant__ int d_nDraws

Definition at line 9 of file gpuMCMCProcessorUtils.cu.

d_nEntries

__device__ __constant__ int d_nEntries

Definition at line 10 of file gpuMCMCProcessorUtils.cu.

d_nLag

__device__ __constant__ int d_nLag

Definition at line 8 of file gpuMCMCProcessorUtils.cu.