Commit c1921f0d authored by BRAMAS Berenger's avatar BRAMAS Berenger
Browse files

Should do P2P until counter in opencl

parent ddf35b57
...@@ -473,7 +473,7 @@ void P2P(const int3 pos, ...@@ -473,7 +473,7 @@ void P2P(const int3 pos,
void P2PRemote(const int3 pos, void P2PRemote(const int3 pos,
struct FOpenCLGroupAttachedLeaf targets, const struct FOpenCLGroupAttachedLeaf sources, struct FOpenCLGroupAttachedLeaf targets, const struct FOpenCLGroupAttachedLeaf sources,
struct FOpenCLGroupAttachedLeaf directNeighborsParticles, const int position, __global void* user_data){ struct FOpenCLGroupAttachedLeaf directNeighborsParticles, const int position, __global void* user_data){
return;// TODO //return;// TODO
__global long long* partdown = targets.attributes[0]; __global long long* partdown = targets.attributes[0];
for(int idxPart = 0 ; idxPart < targets.nbParticles ; ++idxPart){ for(int idxPart = 0 ; idxPart < targets.nbParticles ; ++idxPart){
partdown[idxPart] += directNeighborsParticles.nbParticles; partdown[idxPart] += directNeighborsParticles.nbParticles;
...@@ -503,7 +503,7 @@ int3 getCoordinate(__global const unsigned char* cell, __global void* user_data) ...@@ -503,7 +503,7 @@ int3 getCoordinate(__global const unsigned char* cell, __global void* user_data)
#define FOpenCLCheck( test ) { FOpenCLCheckCore((test), __FILE__, __LINE__); } #define FOpenCLCheck( test ) { FOpenCLCheckCore((test), __FILE__, __LINE__); }
#define FOpenCLCheckAfterCall() { FOpenCLCheckCore((cudaGetLastError()), __FILE__, __LINE__); } #define FOpenCLCheckAfterCall() { FOpenCLCheckCore((cudaGetLastError()), __FILE__, __LINE__); }
#define FOpenCLAssertLF(ARGS) if(!(ARGS)){} #define FOpenCLAssertLF(ARGS) if(!(ARGS)){ *((char*)0x09) = 'e'; }
//#define FOpenCLAssertLF(ARGS) ARGS; //#define FOpenCLAssertLF(ARGS) ARGS;
#define FMGetOppositeNeighIndex(index) (27-(index)-1) #define FMGetOppositeNeighIndex(index) (27-(index)-1)
...@@ -831,7 +831,7 @@ __kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containers ...@@ -831,7 +831,7 @@ __kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containers
struct FOpenCLGroupAttachedLeaf particles = FOpenCLGroupOfParticles_getLeaf(&containers, outsideInteractions[outInterIdx].insideIndex); struct FOpenCLGroupAttachedLeaf particles = FOpenCLGroupOfParticles_getLeaf(&containers, outsideInteractions[outInterIdx].insideIndex);
FOpenCLAssertLF(FOpenCLGroupAttachedLeaf_isAttachedToSomething(&particles)); FOpenCLAssertLF(FOpenCLGroupAttachedLeaf_isAttachedToSomething(&particles));
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].outPosition, 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].outPosition), 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