Commit c5cf75a3 authored by BRAMAS Berenger's avatar BRAMAS Berenger

Update starpu fmm

parent a7996fd1
......@@ -175,7 +175,7 @@ protected:
OutOfBlockInteraction property;
property.insideIndex = mindex;
property.outIndex = interactionsIndexes[idxInter];
property.outPosition = interactionsPosition[idxInter];
property.relativeOutPosition = interactionsPosition[idxInter];
property.insideIdxInBlock = cellIdx;
outsideInteractions.push_back(property);
}
......@@ -215,8 +215,8 @@ protected:
FAssertLF(cell.getMortonIndex() == outsideInteractions[outInterIdx].insideIndex);
const CellClass* ptCell = &interCell;
kernels->M2L( &cell , &ptCell, &outsideInteractions[outInterIdx].outPosition, 1, idxLevel);
const int otherPos = getOppositeInterIndex(outsideInteractions[outInterIdx].outPosition);
kernels->M2L( &cell , &ptCell, &outsideInteractions[outInterIdx].relativeOutPosition, 1, idxLevel);
const int otherPos = getOppositeInterIndex(outsideInteractions[outInterIdx].relativeOutPosition);
ptCell = &cell;
kernels->M2L( &interCell , &ptCell, &otherPos, 1, idxLevel);
}
......@@ -345,7 +345,7 @@ protected:
OutOfBlockInteraction property;
property.insideIndex = mindex;
property.outIndex = interactionsIndexes[idxInter];
property.outPosition = interactionsPosition[idxInter];
property.relativeOutPosition = interactionsPosition[idxInter];
property.insideIdxInBlock = leafIdx;
outsideInteractions.push_back(property);
}
......@@ -386,8 +386,8 @@ protected:
ParticleContainerClass* ptrLeaf = &interParticles;
kernels->P2POuter( FTreeCoordinate(outsideInteractions[outInterIdx].insideIndex, tree->getHeight()-1),
&particles , &ptrLeaf, &outsideInteractions[outInterIdx].outPosition, 1);
const int otherPosition = getOppositeNeighIndex(outsideInteractions[outInterIdx].outPosition);
&particles , &ptrLeaf, &outsideInteractions[outInterIdx].relativeOutPosition, 1);
const int otherPosition = getOppositeNeighIndex(outsideInteractions[outInterIdx].relativeOutPosition);
ptrLeaf = &particles;
kernels->P2POuter( FTreeCoordinate(outsideInteractions[outInterIdx].outIndex, tree->getHeight()-1),
&interParticles , &ptrLeaf, &otherPosition, 1);
......
This diff is collapsed.
......@@ -39,6 +39,8 @@
#endif
#include "Containers/FBoolArray.hpp"
template <class OctreeClass, class CellContainerClass, class KernelClass, class ParticleGroupClass, class StarPUCpuWrapperClass
#ifdef SCALFMM_ENABLE_CUDA_KERNEL
, class StarPUCudaWrapperClass = FStarPUCudaWrapper<KernelClass, FCudaEmptyCellSymb, int, int, FCudaGroupOfCells<FCudaEmptyCellSymb, int, int>,
......@@ -624,8 +626,9 @@ protected:
OutOfBlockInteraction property;
property.insideIndex = mindex;
property.outIndex = interactionsIndexes[idxInter];
property.outPosition = interactionsPosition[idxInter];
property.relativeOutPosition = interactionsPosition[idxInter];
property.insideIdxInBlock = leafIdx;
property.outsideIdxInBlock = -1;
outsideInteractions.push_back(property);
}
}
......@@ -640,16 +643,28 @@ protected:
const MortonIndex blockStartIdxOther = leftContainers->getStartingIndex();
const MortonIndex blockEndIdxOther = leftContainers->getEndingIndex();
while(currentOutInteraction < int(outsideInteractions.size()) && outsideInteractions[currentOutInteraction].outIndex < blockStartIdxOther){
while(currentOutInteraction < int(outsideInteractions.size())
&& (outsideInteractions[currentOutInteraction].outIndex < blockStartIdxOther
|| leftContainers->getLeafIndex(outsideInteractions[currentOutInteraction].outIndex) == -1)
&& outsideInteractions[currentOutInteraction].outIndex < blockEndIdxOther){
currentOutInteraction += 1;
}
int lastOutInteraction = currentOutInteraction;
int copyExistingInteraction = currentOutInteraction;
while(lastOutInteraction < int(outsideInteractions.size()) && outsideInteractions[lastOutInteraction].outIndex < blockEndIdxOther){
const int leafPos = leftContainers->getLeafIndex(outsideInteractions[lastOutInteraction].outIndex);
if(leafPos != -1){
if(copyExistingInteraction != lastOutInteraction){
outsideInteractions[copyExistingInteraction] = outsideInteractions[lastOutInteraction];
}
outsideInteractions[copyExistingInteraction].outsideIdxInBlock = leafPos;
copyExistingInteraction += 1;
}
lastOutInteraction += 1;
}
const int nbInteractionsBetweenBlocks = (lastOutInteraction-currentOutInteraction);
const int nbInteractionsBetweenBlocks = (copyExistingInteraction-currentOutInteraction);
if(nbInteractionsBetweenBlocks){
externalInteractions->emplace_back();
BlockInteractions<ParticleGroupClass>* interactions = &externalInteractions->back();
......@@ -657,7 +672,7 @@ protected:
interactions->otherBlockId = idxLeftGroup;
interactions->interactions.resize(nbInteractionsBetweenBlocks);
std::copy(outsideInteractions.begin() + currentOutInteraction,
outsideInteractions.begin() + lastOutInteraction,
outsideInteractions.begin() + copyExistingInteraction,
interactions->interactions.begin());
}
......@@ -699,8 +714,9 @@ protected:
OutOfBlockInteraction property;
property.insideIndex = mindex;
property.outIndex = interactionsIndexes[idxInter];
property.outPosition = interactionsPosition[idxInter];
property.relativeOutPosition = interactionsPosition[idxInter];
property.insideIdxInBlock = cellIdx;
property.outsideIdxInBlock = -1;
outsideInteractions.push_back(property);
}
}
......@@ -715,17 +731,29 @@ protected:
const MortonIndex blockStartIdxOther = leftCells->getStartingIndex();
const MortonIndex blockEndIdxOther = leftCells->getEndingIndex();
while(currentOutInteraction < int(outsideInteractions.size()) && outsideInteractions[currentOutInteraction].outIndex < blockStartIdxOther){
while(currentOutInteraction < int(outsideInteractions.size())
&& (outsideInteractions[currentOutInteraction].outIndex < blockStartIdxOther
|| leftCells->getCellIndex(outsideInteractions[currentOutInteraction].outIndex) == -1)
&& outsideInteractions[currentOutInteraction].outIndex < blockEndIdxOther){
currentOutInteraction += 1;
}
int lastOutInteraction = currentOutInteraction;
int copyExistingInteraction = currentOutInteraction;
while(lastOutInteraction < int(outsideInteractions.size()) && outsideInteractions[lastOutInteraction].outIndex < blockEndIdxOther){
const int cellPos = leftCells->getCellIndex(outsideInteractions[lastOutInteraction].outIndex);
if(cellPos != -1){
if(copyExistingInteraction != lastOutInteraction){
outsideInteractions[copyExistingInteraction] = outsideInteractions[lastOutInteraction];
}
outsideInteractions[copyExistingInteraction].outsideIdxInBlock = cellPos;
copyExistingInteraction += 1;
}
lastOutInteraction += 1;
}
// Create interactions
const int nbInteractionsBetweenBlocks = (lastOutInteraction-currentOutInteraction);
const int nbInteractionsBetweenBlocks = (copyExistingInteraction-currentOutInteraction);
if(nbInteractionsBetweenBlocks){
externalInteractions->emplace_back();
BlockInteractions<CellContainerClass>* interactions = &externalInteractions->back();
......@@ -733,7 +761,7 @@ protected:
interactions->otherBlockId = idxLeftGroup;
interactions->interactions.resize(nbInteractionsBetweenBlocks);
std::copy(outsideInteractions.begin() + currentOutInteraction,
outsideInteractions.begin() + lastOutInteraction,
outsideInteractions.begin() + copyExistingInteraction,
interactions->interactions.begin());
}
......
......@@ -9,8 +9,9 @@
struct alignas(FStarPUDefaultAlign::StructAlign) OutOfBlockInteraction{
MortonIndex outIndex;
MortonIndex insideIndex;
int outPosition;
int relativeOutPosition;
int insideIdxInBlock;
int outsideIdxInBlock;
// To sort
bool operator <=(const OutOfBlockInteraction& other) const{
return outIndex <= other.outIndex;
......
......@@ -171,7 +171,7 @@ __global__ void FCuda__transferInoutPassPerformMpi(unsigned char* currentCellsP
typename CellContainerClass::CompleteCellClass interactions[343];
memset(interactions, 0, 343*sizeof(interactions[0]));
interactions[outsideInteractions[outInterIdx].outPosition] = interCell;
interactions[outsideInteractions[outInterIdx].relativeOutPosition] = interCell;
const int counter = 1;
kernel->M2L( cell , interactions, counter, idxLevel);
}
......@@ -295,12 +295,12 @@ __global__ void FCuda__transferInoutPassPerform(unsigned char* currentCellsPtr,
typename CellContainerClass::CompleteCellClass interactions[343];
memset(interactions, 0, 343*sizeof(interactions[0]));
interactions[outsideInteractions[outInterIdx].outPosition] = interCell;
interactions[outsideInteractions[outInterIdx].relativeOutPosition] = interCell;
const int counter = 1;
kernel->M2L( cell , interactions, counter, idxLevel);
interactions[outsideInteractions[outInterIdx].outPosition].symb = nullptr;
interactions[FMGetOppositeInterIndex(outsideInteractions[outInterIdx].outPosition)] = cell;
interactions[outsideInteractions[outInterIdx].relativeOutPosition].symb = nullptr;
interactions[FMGetOppositeInterIndex(outsideInteractions[outInterIdx].relativeOutPosition)] = cell;
kernel->M2L( interCell , interactions, counter, idxLevel);
}
}
......@@ -444,7 +444,7 @@ __global__ void FCuda__directInoutPassPerformMpi(unsigned char* containersPtr, s
ParticleGroupClass particles = containers.template getLeaf<ParticleGroupClass>(outsideInteractions[outInterIdx].insideIdxInBlock);
ParticleGroupClass* interactions[27];
memset(interactions, 0, 27*sizeof(ParticleGroupClass*));
interactions[outsideInteractions[outInterIdx].outPosition] = &interParticles;
interactions[outsideInteractions[outInterIdx].relativeOutPosition] = &interParticles;
const int counter = 1;
kernel->P2PRemote( FCudaTreeCoordinate::GetPositionFromMorton(outsideInteractions[outInterIdx].insideIndex, treeHeight-1), &particles, &particles , interactions, counter);
}
......@@ -561,12 +561,12 @@ __global__ void FCuda__directInoutPassPerform(unsigned char* containersPtr, std:
ParticleGroupClass* interactions[27];
memset(interactions, 0, 27*sizeof(ParticleGroupClass*));
interactions[outsideInteractions[outInterIdx].outPosition] = &interParticles;
interactions[outsideInteractions[outInterIdx].relativeOutPosition] = &interParticles;
const int counter = 1;
kernel->P2PRemote( FCudaTreeCoordinate::GetPositionFromMorton(outsideInteractions[outInterIdx].insideIndex, treeHeight-1), &particles, &particles , interactions, counter);
interactions[outsideInteractions[outInterIdx].outPosition] = nullptr;
interactions[FMGetOppositeNeighIndex(outsideInteractions[outInterIdx].outPosition)] = &particles;
interactions[outsideInteractions[outInterIdx].relativeOutPosition] = nullptr;
interactions[FMGetOppositeNeighIndex(outsideInteractions[outInterIdx].relativeOutPosition)] = &particles;
kernel->P2PRemote( FCudaTreeCoordinate::GetPositionFromMorton(outsideInteractions[outInterIdx].outIndex, treeHeight-1), &interParticles, &interParticles , interactions, counter);
}
}
......
......@@ -8,7 +8,7 @@ typedef long long int MortonIndex;
struct OutOfBlockInteraction{
MortonIndex outIndex;
MortonIndex insideIndex;
int outPosition;
int relativeOutPosition;
} __attribute__ ((aligned (DefaultStructAlign)));
struct Uptr9{
__global unsigned char* ptrs[9];
......
......@@ -20,7 +20,7 @@ public:
struct OutOfBlockInteraction{\
MortonIndex outIndex;\
MortonIndex insideIndex;\
int outPosition;\
int relativeOutPosition;\
int insideIdxInBlock;\
} __attribute__ ((aligned (DefaultStructAlign)));\
struct Uptr9{\
......
......@@ -204,7 +204,7 @@ public:
CellClass cell = currentCells->getDownCell((*outsideInteractions)[outInterIdx].insideIdxInBlock);
FAssertLF(cell.getMortonIndex() == (*outsideInteractions)[outInterIdx].insideIndex);
const CellClass* ptCell = &interCell;
kernel->M2L( &cell , &ptCell, &(*outsideInteractions)[outInterIdx].outPosition, 1, idxLevel);
kernel->M2L( &cell , &ptCell, &(*outsideInteractions)[outInterIdx].relativeOutPosition, 1, idxLevel);
}
}
}
......@@ -294,31 +294,25 @@ public:
if(mode == 1){
for(int outInterIdx = 0 ; outInterIdx < int(outsideInteractions->size()) ; ++outInterIdx){
const int cellPos = cellsOther->getCellIndex((*outsideInteractions)[outInterIdx].outIndex);
if(cellPos != -1){
CellClass interCell = cellsOther->getUpCell(cellPos);
FAssertLF(interCell.getMortonIndex() == (*outsideInteractions)[outInterIdx].outIndex);
CellClass cell = currentCells->getDownCell((*outsideInteractions)[outInterIdx].insideIdxInBlock);
FAssertLF(cell.getMortonIndex() == (*outsideInteractions)[outInterIdx].insideIndex);
const CellClass* ptCell = &interCell;
kernel->M2L( &cell , &ptCell, &(*outsideInteractions)[outInterIdx].outPosition, 1, idxLevel);
}
CellClass interCell = cellsOther->getUpCell((*outsideInteractions)[outInterIdx].outsideIdxInBlock);
FAssertLF(interCell.getMortonIndex() == (*outsideInteractions)[outInterIdx].outIndex);
CellClass cell = currentCells->getDownCell((*outsideInteractions)[outInterIdx].insideIdxInBlock);
FAssertLF(cell.getMortonIndex() == (*outsideInteractions)[outInterIdx].insideIndex);
const CellClass* ptCell = &interCell;
kernel->M2L( &cell , &ptCell, &(*outsideInteractions)[outInterIdx].relativeOutPosition, 1, idxLevel);
}
}
else{
for(int outInterIdx = 0 ; outInterIdx < int(outsideInteractions->size()) ; ++outInterIdx){
const int cellPos = currentCells->getCellIndex((*outsideInteractions)[outInterIdx].outIndex);
if(cellPos != -1){
CellClass cell = cellsOther->getUpCell((*outsideInteractions)[outInterIdx].insideIdxInBlock);
FAssertLF(cell.getMortonIndex() == (*outsideInteractions)[outInterIdx].insideIndex);
CellClass interCell = currentCells->getDownCell(cellPos);
FAssertLF(interCell.getMortonIndex() == (*outsideInteractions)[outInterIdx].outIndex);
const int otherPos = getOppositeInterIndex((*outsideInteractions)[outInterIdx].outPosition);
const CellClass* ptCell = &cell;
kernel->M2L( &interCell , &ptCell, &otherPos, 1, idxLevel);
}
CellClass cell = cellsOther->getUpCell((*outsideInteractions)[outInterIdx].insideIdxInBlock);
FAssertLF(cell.getMortonIndex() == (*outsideInteractions)[outInterIdx].insideIndex);
CellClass interCell = currentCells->getDownCell((*outsideInteractions)[outInterIdx].outsideIdxInBlock);
FAssertLF(interCell.getMortonIndex() == (*outsideInteractions)[outInterIdx].outIndex);
const int otherPos = getOppositeInterIndex((*outsideInteractions)[outInterIdx].relativeOutPosition);
const CellClass* ptCell = &cell;
kernel->M2L( &interCell , &ptCell, &otherPos, 1, idxLevel);
}
}
}
......@@ -421,7 +415,7 @@ public:
FAssertLF(containers->getLeafMortonIndex(leafPos) == (*outsideInteractions)[outInterIdx].insideIndex);
ParticleContainerClass* ptrLeaf = &interParticles;
kernel->P2PRemote( FTreeCoordinate((*outsideInteractions)[outInterIdx].insideIndex, treeHeight-1), &particles, &particles ,
&ptrLeaf, &(*outsideInteractions)[outInterIdx].outPosition, 1);
&ptrLeaf, &(*outsideInteractions)[outInterIdx].relativeOutPosition, 1);
}
}
}
......@@ -494,22 +488,19 @@ public:
const std::vector<OutOfBlockInteraction>* outsideInteractions){
KernelClass*const kernel = kernels[omp_get_thread_num()];
for(int outInterIdx = 0 ; outInterIdx < int(outsideInteractions->size()) ; ++outInterIdx){
const int leafPos = containersOther->getLeafIndex((*outsideInteractions)[outInterIdx].outIndex);
if(leafPos != -1){
ParticleContainerClass interParticles = containersOther->template getLeaf<ParticleContainerClass>(leafPos);
ParticleContainerClass particles = containers->template getLeaf<ParticleContainerClass>((*outsideInteractions)[outInterIdx].insideIdxInBlock);
FAssertLF(containersOther->getLeafMortonIndex(leafPos) == (*outsideInteractions)[outInterIdx].outIndex);
FAssertLF(containers->getLeafMortonIndex((*outsideInteractions)[outInterIdx].insideIdxInBlock) == (*outsideInteractions)[outInterIdx].insideIndex);
ParticleContainerClass* ptrLeaf = &interParticles;
kernel->P2POuter( FTreeCoordinate((*outsideInteractions)[outInterIdx].insideIndex, treeHeight-1),
&particles , &ptrLeaf, &(*outsideInteractions)[outInterIdx].outPosition, 1);
const int otherPosition = getOppositeNeighIndex((*outsideInteractions)[outInterIdx].outPosition);
ptrLeaf = &particles;
kernel->P2POuter( FTreeCoordinate((*outsideInteractions)[outInterIdx].outIndex, treeHeight-1),
ParticleContainerClass interParticles = containersOther->template getLeaf<ParticleContainerClass>((*outsideInteractions)[outInterIdx].outsideIdxInBlock);
ParticleContainerClass particles = containers->template getLeaf<ParticleContainerClass>((*outsideInteractions)[outInterIdx].insideIdxInBlock);
FAssertLF(containersOther->getLeafMortonIndex((*outsideInteractions)[outInterIdx].outsideIdxInBlock) == (*outsideInteractions)[outInterIdx].outIndex);
FAssertLF(containers->getLeafMortonIndex((*outsideInteractions)[outInterIdx].insideIdxInBlock) == (*outsideInteractions)[outInterIdx].insideIndex);
ParticleContainerClass* ptrLeaf = &interParticles;
kernel->P2POuter( FTreeCoordinate((*outsideInteractions)[outInterIdx].insideIndex, treeHeight-1),
&particles , &ptrLeaf, &(*outsideInteractions)[outInterIdx].relativeOutPosition, 1);
const int otherPosition = getOppositeNeighIndex((*outsideInteractions)[outInterIdx].relativeOutPosition);
ptrLeaf = &particles;
kernel->P2POuter( FTreeCoordinate((*outsideInteractions)[outInterIdx].outIndex, treeHeight-1),
&interParticles , &ptrLeaf, &otherPosition, 1);
}
}
}
......
......@@ -49,7 +49,7 @@ struct FWrappeCell{
struct OutOfBlockInteraction{
MortonIndex outIndex;
MortonIndex insideIndex;
int outPosition;
int relativeOutPosition;
int insideIdxInBlock;
} __attribute__ ((aligned (DefaultStructAlign)));
......@@ -725,7 +725,7 @@ __kernel void FOpenCL__transferInoutPassPerformMpi(__global unsigned char* curr
struct FWrappeCell interactions[343];
FSetToNullptr343(interactions);
interactions[outsideInteractions[outInterIdx].outPosition] = interCell;
interactions[outsideInteractions[outInterIdx].relativeOutPosition] = interCell;
const int counter = 1;
M2L( cell , interactions, counter, idxLevel, userkernel);
}
......@@ -800,12 +800,12 @@ __kernel void FOpenCL__transferInoutPassPerform(__global unsigned char* currentC
struct FWrappeCell interactions[343];
FSetToNullptr343(interactions);
interactions[outsideInteractions[outInterIdx].outPosition] = interCell;
interactions[outsideInteractions[outInterIdx].relativeOutPosition] = interCell;
const int counter = 1;
M2L( cell , interactions, counter, idxLevel, userkernel);
interactions[outsideInteractions[outInterIdx].outPosition].symb = NULLPTR;
interactions[FMGetOppositeInterIndex(outsideInteractions[outInterIdx].outPosition)] = cell;
interactions[outsideInteractions[outInterIdx].relativeOutPosition].symb = NULLPTR;
interactions[FMGetOppositeInterIndex(outsideInteractions[outInterIdx].relativeOutPosition)] = cell;
M2L( interCell , interactions, counter, idxLevel, userkernel);
}
}
......@@ -884,7 +884,7 @@ __kernel void FOpenCL__directInoutPassPerformMpi(__global unsigned char* contain
struct FOpenCLGroupAttachedLeaf particles = FOpenCLGroupOfParticles_getLeaf(&containers, outsideInteractions[outInterIdx].insideIdxInBlock);
FOpenCLAssertLF(FOpenCLGroupOfParticles_getLeafMortonIndex(&containers, outsideInteractions[outInterIdx].insideIdxInBlock) == outsideInteractions[outInterIdx].insideIndex);
P2PRemote( GetPositionFromMorton(outsideInteractions[outInterIdx].insideIndex, treeHeight-1), particles, particles , interParticles, outsideInteractions[outInterIdx].outPosition, userkernel);
P2PRemote( GetPositionFromMorton(outsideInteractions[outInterIdx].insideIndex, treeHeight-1), particles, particles , interParticles, outsideInteractions[outInterIdx].relativeOutPosition, userkernel);
}
}
}
......@@ -951,9 +951,9 @@ __kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containers
FOpenCLAssertLF(particles.nbParticles);
FOpenCLAssertLF(interParticles.nbParticles);
P2PRemote( GetPositionFromMorton(outsideInteractions[outInterIdx].insideIndex, treeHeight-1), particles, particles , interParticles, outsideInteractions[outInterIdx].outPosition, userkernel );
P2PRemote( GetPositionFromMorton(outsideInteractions[outInterIdx].insideIndex, treeHeight-1), particles, particles , interParticles, outsideInteractions[outInterIdx].relativeOutPosition, userkernel );
P2PRemote( GetPositionFromMorton(outsideInteractions[outInterIdx].outIndex, treeHeight-1), interParticles, interParticles , particles, FMGetOppositeNeighIndex(outsideInteractions[outInterIdx].outPosition), userkernel);
P2PRemote( GetPositionFromMorton(outsideInteractions[outInterIdx].outIndex, treeHeight-1), interParticles, interParticles , particles, FMGetOppositeNeighIndex(outsideInteractions[outInterIdx].relativeOutPosition), userkernel);
}
}
}
......
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment