Commit 725ebabf authored by BRAMAS Berenger's avatar BRAMAS Berenger

Update blocked tree to have lighter blocks - the tables of indexes has been...

Update blocked tree to have lighter blocks - the tables of indexes has been removed for non existing cells - should be tested and the opencl code is not up to date for now
parent e6e5ca8f
This diff is collapsed.
This diff is collapsed.
......@@ -25,7 +25,6 @@ class FGroupOfParticles {
MortonIndex startingIndex;
MortonIndex endingIndex;
int numberOfLeavesInBlock;
int blockIndexesTableSize;
//< The real number of particles allocated
FSize nbParticlesAllocatedInGroup;
......@@ -41,6 +40,7 @@ class FGroupOfParticles {
/** Information about a leaf */
struct alignas(FStarPUDefaultAlign::StructAlign) LeafHeader {
MortonIndex mindex;
FSize nbParticles;
size_t offSet;
};
......@@ -60,9 +60,6 @@ protected:
return nbParticles + (MemoryAlignementParticles - (nbParticles%MemoryAlignementParticles));
}
//< This value is for not used leaves
static const int LeafIsEmptyFlag = -1;
//< The size of memoryBuffer in byte
size_t allocatedMemoryInByte;
//< Pointer to a block memory
......@@ -70,8 +67,6 @@ protected:
//< Pointer to the header inside the block memory
BlockHeader* blockHeader;
//< Pointer to the indexes table inside the block memory
int* blockIndexesTable;
//< Pointer to leaves information
LeafHeader* leafHeader;
//< The total number of particles in the group
......@@ -96,13 +91,11 @@ public:
FGroupOfParticles(unsigned char* inBuffer, const size_t inAllocatedMemoryInByte,
unsigned char* inAttributes)
: allocatedMemoryInByte(inAllocatedMemoryInByte), memoryBuffer(inBuffer),
blockHeader(nullptr), blockIndexesTable(nullptr), leafHeader(nullptr), nbParticlesInGroup(0),
blockHeader(nullptr), leafHeader(nullptr), nbParticlesInGroup(0),
attributesBuffer(nullptr), deleteBuffer(false){
// Move the pointers to the correct position
blockHeader = reinterpret_cast<BlockHeader*>(inBuffer);
inBuffer += sizeof(BlockHeader);
blockIndexesTable = reinterpret_cast<int*>(inBuffer);
inBuffer += (blockHeader->blockIndexesTableSize*sizeof(int));
leafHeader = reinterpret_cast<LeafHeader*>(inBuffer);
// Init particle pointers
......@@ -133,7 +126,7 @@ public:
* @param inNumberOfLeaves total number of leaves in the interval (should be <= inEndingIndex-inEndingIndex)
*/
FGroupOfParticles(const MortonIndex inStartingIndex, const MortonIndex inEndingIndex, const int inNumberOfLeaves, const FSize inNbParticles)
: allocatedMemoryInByte(0), memoryBuffer(nullptr), blockHeader(nullptr), blockIndexesTable(nullptr), leafHeader(nullptr), nbParticlesInGroup(inNbParticles),
: allocatedMemoryInByte(0), memoryBuffer(nullptr), blockHeader(nullptr), leafHeader(nullptr), nbParticlesInGroup(inNbParticles),
deleteBuffer(true){
memset(particlePosition, 0, sizeof(particlePosition));
memset(particleAttributes, 0, sizeof(particleAttributes));
......@@ -141,12 +134,11 @@ public:
const FSize nbParticlesAllocatedInGroup = RoundToUpperParticles(nbParticlesInGroup+(MemoryAlignementParticles-1)*inNumberOfLeaves);
// Find the number of leaf to allocate in the blocks
const int blockIndexesTableSize = int(inEndingIndex-inStartingIndex);
FAssertLF(inNumberOfLeaves <= blockIndexesTableSize);
FAssertLF(int(inEndingIndex-inStartingIndex) >= inNumberOfLeaves);
// Total number of bytes in the block
const size_t sizeOfOneParticle = (3*sizeof(FReal) + NbSymbAttributes*sizeof(AttributeClass));
const size_t memoryToAlloc = sizeof(BlockHeader)
+ (blockIndexesTableSize*sizeof(int))
+ (inNumberOfLeaves*sizeof(MortonIndex))
+ (inNumberOfLeaves*sizeof(LeafHeader))
+ nbParticlesAllocatedInGroup*sizeOfOneParticle;
......@@ -161,15 +153,12 @@ public:
unsigned char* bufferPtr = memoryBuffer;
blockHeader = reinterpret_cast<BlockHeader*>(bufferPtr);
bufferPtr += sizeof(BlockHeader);
blockIndexesTable = reinterpret_cast<int*>(bufferPtr);
bufferPtr += (blockIndexesTableSize*sizeof(int));
leafHeader = reinterpret_cast<LeafHeader*>(bufferPtr);
// Init header
blockHeader->startingIndex = inStartingIndex;
blockHeader->endingIndex = inEndingIndex;
blockHeader->numberOfLeavesInBlock = inNumberOfLeaves;
blockHeader->blockIndexesTableSize = blockIndexesTableSize;
blockHeader->nbParticlesAllocatedInGroup = nbParticlesAllocatedInGroup;
// Init particle pointers
......@@ -197,8 +186,8 @@ public:
}
// Set all index to not used
for(int idxLeafPtr = 0 ; idxLeafPtr < blockIndexesTableSize ; ++idxLeafPtr){
blockIndexesTable[idxLeafPtr] = LeafIsEmptyFlag;
for(int idxLeafPtr = 0 ; idxLeafPtr < inNumberOfLeaves ; ++idxLeafPtr){
leafHeader[idxLeafPtr].mindex = -1;
}
}
......@@ -262,7 +251,7 @@ public:
/** The size of the interval endingIndex-startingIndex (set from the constructor) */
int getSizeOfInterval() const {
return blockHeader->blockIndexesTableSize;
return int(blockHeader->endingIndex-blockHeader->startingIndex);
}
/** Return true if inIndex should be located in the current block */
......@@ -270,20 +259,45 @@ public:
return blockHeader->startingIndex <= inIndex && inIndex < blockHeader->endingIndex;
}
/** Return the idx in array of the cell */
MortonIndex getLeafMortonIndex(const int id) const{
FAssertLF(id < blockHeader->numberOfLeavesInBlock);
return leafHeader[id].mindex;
}
/** Check if a cell exist (by binary search) and return it index */
int getLeafIndex(const MortonIndex leafIdx) const{
int idxLeft = 0;
int idxRight = blockHeader->numberOfLeavesInBlock-1;
while(idxLeft <= idxRight){
const int idxMiddle = (idxLeft+idxRight)/2;
if(leafHeader[idxMiddle].mindex == leafIdx){
return idxMiddle;
}
if(leafIdx < leafHeader[idxMiddle].mindex){
idxRight = idxMiddle-1;
}
else{
idxLeft = idxMiddle+1;
}
}
return -1;
}
/** Return true if inIndex is located in the current block and is not empty */
bool exists(const MortonIndex inIndex) const {
return isInside(inIndex) && (blockIndexesTable[inIndex-blockHeader->startingIndex] != LeafIsEmptyFlag);
return isInside(inIndex) && (getLeafIndex(inIndex) != -1);
}
/** Allocate a new leaf by calling its constructor */
size_t newLeaf(const MortonIndex inIndex, const int id, const FSize nbParticles, const size_t offsetInGroup){
FAssertLF(isInside(inIndex));
FAssertLF(!exists(inIndex));
FAssertLF(id < blockHeader->blockIndexesTableSize);
FAssertLF(id < blockHeader->numberOfLeavesInBlock);
FAssertLF(offsetInGroup < size_t(blockHeader->nbParticlesAllocatedInGroup));
blockIndexesTable[inIndex-blockHeader->startingIndex] = id;
leafHeader[id].nbParticles = nbParticles;
leafHeader[id].offSet = offsetInGroup;
leafHeader[id].mindex = inIndex;
const size_t nextLeafOffsetInGroup = RoundToUpperParticles(offsetInGroup+nbParticles);
FAssertLF(nextLeafOffsetInGroup <= size_t(blockHeader->nbParticlesAllocatedInGroup + MemoryAlignementParticles));
......@@ -293,32 +307,26 @@ public:
/** Iterate on each allocated leaves */
template<class ParticlesAttachedClass>
void forEachLeaf(std::function<void(ParticlesAttachedClass*)> function){
for(int idxLeafPtr = 0 ; idxLeafPtr < blockHeader->blockIndexesTableSize ; ++idxLeafPtr){
if(blockIndexesTable[idxLeafPtr] != LeafIsEmptyFlag){
const int id = blockIndexesTable[idxLeafPtr];
ParticlesAttachedClass leaf(leafHeader[id].nbParticles,
particlePosition[0] + leafHeader[id].offSet,
for(int idxLeafPtr = 0 ; idxLeafPtr < blockHeader->numberOfLeavesInBlock ; ++idxLeafPtr){
ParticlesAttachedClass leaf(leafHeader[idxLeafPtr].nbParticles,
particlePosition[0] + leafHeader[idxLeafPtr].offSet,
blockHeader->positionsLeadingDim,
(attributesBuffer?particleAttributes[NbSymbAttributes] + leafHeader[id].offSet:nullptr),
(attributesBuffer?particleAttributes[NbSymbAttributes] + leafHeader[idxLeafPtr].offSet:nullptr),
blockHeader->attributeLeadingDim);
function(&leaf);
}
function(&leaf);
}
}
/** Return the address of the leaf if it exists (or NULL) */
template<class ParticlesAttachedClass>
ParticlesAttachedClass getLeaf(const MortonIndex leafIndex){
if(blockIndexesTable[leafIndex - blockHeader->startingIndex] != LeafIsEmptyFlag){
const int id = blockIndexesTable[leafIndex - blockHeader->startingIndex];
return ParticlesAttachedClass(leafHeader[id].nbParticles,
ParticlesAttachedClass getLeaf(const int id){
FAssertLF(id < blockHeader->numberOfLeavesInBlock);
return ParticlesAttachedClass(leafHeader[id].nbParticles,
particlePosition[0] + leafHeader[id].offSet,
blockHeader->positionsLeadingDim,
(attributesBuffer?particleAttributes[NbSymbAttributes] + leafHeader[id].offSet:nullptr),
blockHeader->attributeLeadingDim);
}
return ParticlesAttachedClass();
}
};
......
This diff is collapsed.
This diff is collapsed.
This diff is collapsed.
......@@ -26,9 +26,6 @@ public:
typedef FGroupOfCells<CompositeCellClass, SymboleCellClass, PoleCellClass, LocalCellClass> CellGroupClass;
protected:
//< This value is for not used cells
static const int CellIsEmptyFlag = -1;
//< height of the tree (1 => only the root)
const int treeHeight;
//< max number of cells in a block
......@@ -100,7 +97,7 @@ public:
// Add cell
newBlock->newCell(newNodeIndex, cellIdInBlock);
CompositeCellClass newNode = newBlock->getCompleteCell(newNodeIndex);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(newNodeIndex);
newNode.setCoordinate(newNodeCoordinate);
......@@ -109,7 +106,7 @@ public:
blockIteratorInOctree.getCurrentLeaf()->getSrc()->getNbParticles(),
nbParticlesOffsetBeforeLeaf);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(newNodeIndex);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(cellIdInBlock);
attachedLeaf.copyFromContainer(blockIteratorInOctree.getCurrentLeaf()->getSrc(), 0);
cellIdInBlock += 1;
......@@ -150,7 +147,7 @@ public:
const FTreeCoordinate newNodeCoordinate = blockIteratorInOctree.getCurrentCell()->getCoordinate();
newBlock->newCell(newNodeIndex, cellIdInBlock);
CompositeCellClass newNode = newBlock->getCompleteCell(newNodeIndex);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(newNodeIndex);
newNode.setCoordinate(newNodeCoordinate);
......@@ -261,7 +258,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -271,7 +268,7 @@ public:
nbParticlesOffsetBeforeLeaf = newParticleBlock->newLeaf(currentBlockIndexes[cellIdInBlock], cellIdInBlock,
nbParticlesPerLeaf[cellIdInBlock], nbParticlesOffsetBeforeLeaf);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(currentBlockIndexes[cellIdInBlock]);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(cellIdInBlock);
// Copy each particle from the original position
for(FSize idxPart = 0 ; idxPart < nbParticlesPerLeaf[cellIdInBlock] ; ++idxPart){
attachedLeaf.setParticle(idxPart, particlesToSort[idxPart + offsetParticles].originalIndex, inParticlesContainer);
......@@ -345,7 +342,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -467,7 +464,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -477,7 +474,7 @@ public:
nbParticlesOffsetBeforeLeaf = newParticleBlock->newLeaf(currentBlockIndexes[cellIdInBlock], cellIdInBlock,
nbParticlesPerLeaf[cellIdInBlock], nbParticlesOffsetBeforeLeaf);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(currentBlockIndexes[cellIdInBlock]);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(cellIdInBlock);
// Copy each particle from the original position
for(FSize idxPart = 0 ; idxPart < nbParticlesPerLeaf[cellIdInBlock] ; ++idxPart){
attachedLeaf.setParticle(idxPart, particlesToSort[idxPart + offsetParticles].originalIndex, inParticlesContainer);
......@@ -552,7 +549,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -607,7 +604,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -697,7 +694,9 @@ public:
while(iterCells != iterEndCells && iterLeaves != iterEndLeaves){
(*iterCells)->forEachCell([&](CompositeCellClass aCell){
ParticlesAttachedClass aLeaf = (*iterLeaves)->template getLeaf <ParticlesAttachedClass>(aCell.getMortonIndex());
const int leafIdx = (*iterLeaves)->getLeafIndex(aCell.getMortonIndex());
FAssertLF(leafIdx != -1);
ParticlesAttachedClass aLeaf = (*iterLeaves)->template getLeaf <ParticlesAttachedClass>(leafIdx);
FAssertLF(aLeaf.isAttachedToSomething());
function(aCell, &aLeaf);
});
......
......@@ -27,9 +27,6 @@ public:
typedef FGroupOfCellsDyn<CompositeCellClass> CellGroupClass;
protected:
//< This value is for not used cells
static const int CellIsEmptyFlag = -1;
//< height of the tree (1 => only the root)
const int treeHeight;
//< max number of cells in a block
......@@ -109,7 +106,7 @@ public:
// Add cell
newBlock->newCell(newNodeIndex, cellIdInBlock, BuildCellFunc, idxLevel);
CompositeCellClass newNode = newBlock->getCompleteCell(newNodeIndex);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(newNodeIndex);
newNode.setCoordinate(newNodeCoordinate);
......@@ -146,7 +143,7 @@ public:
// Add leaf
newParticleBlock->newLeaf(newNodeIndex, cellIdInBlock);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(newNodeIndex);
BasicAttachedClass attachedLeaf = newParticleBlock->template getLeaf<BasicAttachedClass>(cellIdInBlock);
attachedLeaf.copyFromContainer(newNodeIndex,
blockIteratorLeafInOctree.getCurrentLeaf()->getSrc());
......@@ -189,7 +186,7 @@ public:
const FTreeCoordinate newNodeCoordinate = blockIteratorInOctree.getCurrentCell()->getCoordinate();
newBlock->newCell(newNodeIndex, cellIdInBlock, BuildCellFunc, idxLevel);
CompositeCellClass newNode = newBlock->getCompleteCell(newNodeIndex);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(newNodeIndex);
newNode.setCoordinate(newNodeCoordinate);
......@@ -297,7 +294,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock, BuildCellFunc, idxLevel);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -327,8 +324,8 @@ public:
InitLeafFunc(currentBlockIndexes[cellIdInBlock], &inParticlesContainer[offsetParts],
nbParticlesPerLeaf[cellIdInBlock],
newParticleBlock->getLeafSymbBuffer(currentBlockIndexes[cellIdInBlock]), symbSizePerLeaf[cellIdInBlock],
newParticleBlock->getLeafDownBuffer(currentBlockIndexes[cellIdInBlock]), downSizePerDown[cellIdInBlock]);
newParticleBlock->getLeafSymbBuffer(cellIdInBlock), symbSizePerLeaf[cellIdInBlock],
newParticleBlock->getLeafDownBuffer(cellIdInBlock), downSizePerDown[cellIdInBlock]);
offsetParts += nbParticlesPerLeaf[cellIdInBlock];
}
......@@ -398,7 +395,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock, BuildCellFunc, idxLevel);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -514,7 +511,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock, BuildCellFunc, idxLevel);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -544,8 +541,8 @@ public:
InitLeafFunc(currentBlockIndexes[cellIdInBlock], &inParticlesContainer[offsetParts],
nbParticlesPerLeaf[cellIdInBlock],
newParticleBlock->getLeafSymbBuffer(currentBlockIndexes[cellIdInBlock]), symbSizePerLeaf[cellIdInBlock],
newParticleBlock->getLeafDownBuffer(currentBlockIndexes[cellIdInBlock]), downSizePerDown[cellIdInBlock]);
newParticleBlock->getLeafSymbBuffer(cellIdInBlock), symbSizePerLeaf[cellIdInBlock],
newParticleBlock->getLeafDownBuffer(cellIdInBlock), downSizePerDown[cellIdInBlock]);
offsetParts += nbParticlesPerLeaf[cellIdInBlock];
}
......@@ -616,7 +613,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock, BuildCellFunc, idxLevel);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......@@ -672,7 +669,7 @@ public:
for(int cellIdInBlock = 0; cellIdInBlock != sizeOfBlock ; ++cellIdInBlock){
newBlock->newCell(currentBlockIndexes[cellIdInBlock], cellIdInBlock, BuildCellFunc, idxLevel);
CompositeCellClass newNode = newBlock->getCompleteCell(currentBlockIndexes[cellIdInBlock]);
CompositeCellClass newNode = newBlock->getCompleteCell(cellIdInBlock);
newNode.setMortonIndex(currentBlockIndexes[cellIdInBlock]);
FTreeCoordinate coord;
coord.setPositionFromMorton(currentBlockIndexes[cellIdInBlock], idxLevel);
......
......@@ -10,6 +10,7 @@ struct alignas(FStarPUDefaultAlign::StructAlign) OutOfBlockInteraction{
MortonIndex outIndex;
MortonIndex insideIndex;
int outPosition;
int insideIdxInBlock;
// To sort
bool operator <=(const OutOfBlockInteraction& other) const{
return outIndex <= other.outIndex;
......
This diff is collapsed.
......@@ -16,23 +16,20 @@ class FCudaGroupOfCells {
MortonIndex startingIndex;
MortonIndex endingIndex;
int numberOfCellsInBlock;
int blockIndexesTableSize;
};
protected:
//< The size of the memoryBuffer
int allocatedMemoryInByte;
size_t allocatedMemoryInByte;
//< Pointer to a block memory
unsigned char* memoryBuffer;
//< Pointer to the header inside the block memory
BlockHeader* blockHeader;
//< Pointer to the indexes table inside the block memory
int* blockIndexesTable;
MortonIndex* cellIndexes;
//< Pointer to the cells inside the block memory
SymboleCellClass* blockCells;
//< This value is for not used cells
static const MortonIndex CellIsEmptyFlag = -1;
//< The multipole data
PoleCellClass* cellMultipoles;
......@@ -44,7 +41,7 @@ public:
__device__ FCudaGroupOfCells()
: allocatedMemoryInByte(0), memoryBuffer(nullptr),
blockHeader(nullptr), blockIndexesTable(nullptr), blockCells(nullptr),
blockHeader(nullptr), cellIndexes(nullptr), blockCells(nullptr),
cellMultipoles(nullptr), cellLocals(nullptr){
}
......@@ -55,8 +52,8 @@ public:
memoryBuffer = (inBuffer);
blockHeader = reinterpret_cast<BlockHeader*>(inBuffer);
inBuffer += sizeof(BlockHeader);
blockIndexesTable = reinterpret_cast<int*>(inBuffer);
inBuffer += (blockHeader->blockIndexesTableSize*sizeof(int));
cellIndexes = reinterpret_cast<MortonIndex*>(inBuffer);
inBuffer += (blockHeader->numberOfCellsInBlock*sizeof(MortonIndex));
blockCells = reinterpret_cast<SymboleCellClass*>(inBuffer);
inBuffer += (sizeof(SymboleCellClass)*blockHeader->numberOfCellsInBlock);
//FAssertLF(size_t(inBuffer-memoryBuffer) == allocatedMemoryInByte);
......@@ -73,13 +70,13 @@ public:
__device__ FCudaGroupOfCells(unsigned char* inBuffer, const size_t inAllocatedMemoryInByte,
unsigned char* inCellMultipoles, unsigned char* inCellLocals)
: allocatedMemoryInByte(inAllocatedMemoryInByte), memoryBuffer(inBuffer),
blockHeader(nullptr), blockIndexesTable(nullptr), blockCells(nullptr),
blockHeader(nullptr), cellIndexes(nullptr), blockCells(nullptr),
cellMultipoles(nullptr), cellLocals(nullptr){
// Move the pointers to the correct position
blockHeader = reinterpret_cast<BlockHeader*>(inBuffer);
inBuffer += sizeof(BlockHeader);
blockIndexesTable = reinterpret_cast<int*>(inBuffer);
inBuffer += (blockHeader->blockIndexesTableSize*sizeof(int));
cellIndexes = reinterpret_cast<MortonIndex*>(inBuffer);
inBuffer += (blockHeader->numberOfCellsInBlock*sizeof(MortonIndex));
blockCells = reinterpret_cast<SymboleCellClass*>(inBuffer);
inBuffer += (sizeof(SymboleCellClass)*blockHeader->numberOfCellsInBlock);
//FAssertLF(size_t(inBuffer-memoryBuffer) == allocatedMemoryInByte);
......@@ -104,8 +101,8 @@ public:
}
/** The size of the interval endingIndex-startingIndex (set from the constructor) */
__device__ int getSizeOfInterval() const {
return blockHeader->blockIndexesTableSize;
int getSizeOfInterval() const {
return int(blockHeader->endingIndex-blockHeader->startingIndex);
}
/** Return true if inIndex should be located in the current block */
......@@ -113,69 +110,85 @@ public:
return blockHeader->startingIndex <= inIndex && inIndex < blockHeader->endingIndex;
}
/** Return the idx in array of the cell */
__device__ MortonIndex getCellMortonIndex(const int cellPos) const{
return cellIndexes[cellPos];
}
/** Check if a cell exist (by binary search) and return it index */
__device__ int getFistChildIdx(const MortonIndex parentIdx) const{
int idxLeft = 0;
int idxRight = blockHeader->numberOfCellsInBlock-1;
while(idxLeft <= idxRight){
int idxMiddle = (idxLeft+idxRight)/2;
if((cellIndexes[idxMiddle]>>3) == parentIdx){
while(0 < idxMiddle && (cellIndexes[idxMiddle-1]>>3) == parentIdx){
idxMiddle -= 1;
}
return idxMiddle;
}
if(parentIdx < (cellIndexes[idxMiddle]>>3)){
idxRight = idxMiddle-1;
}
else{
idxLeft = idxMiddle+1;
}
}
return -1;
}
/** Check if a cell exist (by binary search) and return it index */
__device__ int getCellIndex(const MortonIndex cellIdx) const{
int idxLeft = 0;
int idxRight = blockHeader->numberOfCellsInBlock-1;
while(idxLeft <= idxRight){
const int idxMiddle = (idxLeft+idxRight)/2;
if(cellIndexes[idxMiddle] == cellIdx){
return idxMiddle;
}
if(cellIdx < cellIndexes[idxMiddle]){
idxRight = idxMiddle-1;
}
else{
idxLeft = idxMiddle+1;
}
}
return -1;
}
/** Return true if inIndex is located in the current block and is not empty */
__device__ bool exists(const MortonIndex inIndex) const {
return isInside(inIndex) && (blockIndexesTable[inIndex-blockHeader->startingIndex] != CellIsEmptyFlag);
return isInside(inIndex) && (getCellIndex(inIndex) != -1);
}
/** Return the address of the cell if it exists (or NULL) */
__device__ CompleteCellClass getCompleteCell(const MortonIndex inIndex){
__device__ CompleteCellClass getCompleteCell(const int cellPos){
//FAssertLF(cellMultipoles && cellLocals);
if( exists(inIndex) ){
CompleteCellClass cell;
const int cellPos = blockIndexesTable[inIndex-blockHeader->startingIndex];
cell.symb = &blockCells[cellPos];
cell.up = &cellMultipoles[cellPos];
cell.down = &cellLocals[cellPos];
return cell;
}
else{
CompleteCellClass cell;
cell.symb = nullptr;
cell.up = nullptr;
cell.down = nullptr;
return cell;
}
CompleteCellClass cell;
cell.symb = &blockCells[cellPos];
cell.up = &cellMultipoles[cellPos];
cell.down = &cellLocals[cellPos];
return cell;
}
/** Return the address of the cell if it exists (or NULL) */
__device__ CompleteCellClass getUpCell(const MortonIndex inIndex){
__device__ CompleteCellClass getUpCell(const int cellPos){
//FAssertLF(cellMultipoles);
if( exists(inIndex) ){
CompleteCellClass cell;
const int cellPos = blockIndexesTable[inIndex-blockHeader->startingIndex];
cell.symb = &blockCells[cellPos];
cell.up = &cellMultipoles[cellPos];
cell.down = nullptr;
return cell;
}
else{
CompleteCellClass cell;
cell.symb = nullptr;
cell.up = nullptr;
cell.down = nullptr;
return cell;
}
CompleteCellClass cell;
cell.symb = &blockCells[cellPos];
cell.up = &cellMultipoles[cellPos];
cell.down = nullptr;
return cell;
}
/** Return the address of the cell if it exists (or NULL) */
__device__ CompleteCellClass getDownCell(const MortonIndex inIndex){
__device__ CompleteCellClass getDownCell(const int cellPos){
//FAssertLF(cellLocals);
if( exists(inIndex) ){
CompleteCellClass cell;
const int cellPos = blockIndexesTable[inIndex-blockHeader->startingIndex];
cell.symb = &blockCells[cellPos];
cell.up = nullptr;
cell.down = &cellLocals[cellPos];
return cell;
}
else{