__global__ void CountMakroParticle(ParBox parBox, CounterBox counterBox, Mapping mapper)
{
typedef MappingDesc::SuperCellSize SuperCellSize;
typedef typename ParBox::FrameType FrameType;
const DataSpace<simDim> block(mapper.getSuperCellIndex(DataSpace<simDim > (blockIdx)));
/* counterBox has no guarding supercells*/
const DataSpace<simDim> counterCell = block - mapper.getGuardingSuperCells();
const DataSpace<simDim > threadIndex(threadIdx);
const int linearThreadIdx = DataSpaceOperations<simDim>::template map<SuperCellSize > (threadIndex);
__shared__ uint64_cu counterValue;
__shared__ FrameType *frame;
__shared__ bool isValid;
if (linearThreadIdx == 0)
{
counterValue = 0;
frame = &(parBox.getLastFrame(block, isValid));
if (!isValid)
{
counterBox(counterCell) = counterValue;
}
}
__syncthreads();
if (!isValid)
return; //end kernel if we have no frames
bool isParticle = (*frame)[linearThreadIdx][multiMask_];
while (isValid)
{
if (isParticle)
{
atomicAdd(&counterValue, static_cast<uint64_cu> (1LU));
}
__syncthreads();
if (linearThreadIdx == 0)
{
frame = &(parBox.getPreviousFrame(*frame, isValid));
}
isParticle = true;
__syncthreads();
}
if (linearThreadIdx == 0)
counterBox(counterCell) = counterValue;
}
