Commit 1a2c4fdd authored by BRAMAS Berenger's avatar BRAMAS Berenger
Browse files

wip starpu opencl but there still a problem of moving the data back

parent 96763803
......@@ -187,7 +187,6 @@ public:
void upwardPassPerform(cl_mem currentCellsPtr, size_t currentCellsSize, cl_mem subCellGroupsPtr[9], size_t subCellGroupsSize[9], int nbSubCellGroups, int idxLevel){
return; // TODO
Uptr9 ptrs;
memcpy(ptrs.ptrs, subCellGroupsPtr, sizeof(cl_mem)*9);
size_t9 sizes;
......@@ -202,7 +201,6 @@ public:
void transferInoutPassPerformMpi(cl_mem currentCellsPtr,
size_t currentCellsSize, cl_mem externalCellsPtr, size_t externalCellsSize, int idxLevel, cl_mem outsideInteractionsCl,
size_t outsideInteractionsSize){
return; // TODO
SetKernelArgs(kernel_transferInoutPassPerformMpi, 0, &currentCellsPtr,&currentCellsSize, &externalCellsPtr, &externalCellsSize, &idxLevel, &outsideInteractionsCl,
&outsideInteractionsSize, &user_data);
size_t dim = 1;
......@@ -220,7 +218,6 @@ public:
void transferInoutPassPerform(cl_mem currentCellsPtr,
size_t currentCellsSize, cl_mem externalCellsPtr, size_t externalCellsSize, int idxLevel, cl_mem outsideInteractionsCl,
size_t outsideInteractionsSize){
return; // TODO
SetKernelArgs(kernel_transferInoutPassPerform, 0, &currentCellsPtr,&currentCellsSize, &externalCellsPtr, &externalCellsSize, &idxLevel,
&outsideInteractionsCl,&outsideInteractionsSize, &user_data);
size_t dim = 1;
......@@ -230,7 +227,6 @@ public:
void downardPassPerform(cl_mem currentCellsPtr,
size_t currentCellsSize, cl_mem subCellGroupsPtr[9], size_t subCellGroupsSize[9], int nbSubCellGroups, int idxLevel){
return; // TODO
Uptr9 ptrs;
memcpy(ptrs.ptrs, subCellGroupsPtr, sizeof(cl_mem)*9);
size_t9 sizes;
......@@ -246,7 +242,6 @@ public:
void directInoutPassPerformMpi(cl_mem containersPtr,
size_t containersSize, cl_mem externalContainersPtr, size_t externalContainersSize, cl_mem outsideInteractionsCl,
size_t outsideInteractionsSize){
return; // TODO
SetKernelArgs(kernel_directInoutPassPerformMpi, 0, &containersPtr,
&containersSize, &externalContainersPtr, &externalContainersSize, &outsideInteractionsCl,&outsideInteractionsSize, &treeHeight, &user_data);
size_t dim = 1;
......@@ -255,7 +250,6 @@ public:
}
void directInPassPerform(cl_mem containersPtr, size_t containerSize){
return; // TODO
SetKernelArgs(kernel_directInPassPerform, 0, &containersPtr, &containerSize, &treeHeight, &user_data);
size_t dim = 1;
const int err = clEnqueueNDRangeKernel(queue_directInPassPerform, kernel_directInPassPerform, 1, NULL, &dim, NULL, 0, NULL, NULL);
......@@ -265,7 +259,6 @@ public:
void directInoutPassPerform(cl_mem containersPtr,
size_t containerSize, cl_mem externalContainersPtr, size_t externalContainersSize, cl_mem outsideInteractionsCl,
size_t outsideInteractionsSize){
return; // TODO
SetKernelArgs(kernel_directInoutPassPerform, 0, &containersPtr,
&containerSize, &externalContainersPtr, &externalContainersSize, &outsideInteractionsCl, &outsideInteractionsSize, &treeHeight, &user_data);
size_t dim = 1;
......@@ -275,7 +268,6 @@ public:
void mergePassPerform(cl_mem leafCellsPtr,
size_t leafCellsSize, cl_mem containersPtr, size_t containersSize){
return; // TODO
SetKernelArgs(kernel_mergePassPerform, 0, &leafCellsPtr, &leafCellsSize, &containersPtr, &containersSize, &user_data);
size_t dim = 1;
const int err = clEnqueueNDRangeKernel(queue_mergePassPerform, kernel_mergePassPerform, 1, NULL, &dim, NULL, 0, NULL, NULL);
......
......@@ -475,7 +475,10 @@ __kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, s
FOpenCLAssertLF(getMortonIndex(cell, userkernel) == mindex);
struct FOpenCLGroupAttachedLeaf particles = FOpenCLGroupOfParticles_getLeaf(&containers, mindex);
FOpenCLAssertLF(FOpenCLGroupAttachedLeaf_isAttachedToSomething(&particles));
P2M(cell, particles, userkernel);
//P2M(cell, particles, userkernel);
for(int idx = 0 ; idx < FCellClassSize ; ++idx){
cell[idx] = 'z';
}
/*output[0] = blockStartIdx;
output[1] = blockEndIdx;
output[2] = particles.nbParticles;
......@@ -507,7 +510,6 @@ __kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, s
__kernel void FOpenCL__upwardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize,
int nbSubCellGroups, int idxLevel, __global void* userkernel){
return;// TODO
struct FOpenCLGroupOfCells currentCells = BuildFOpenCLGroupOfCells(currentCellsPtr, currentCellsSize);
struct FOpenCLGroupOfCells subCellGroups[9];
for(int idx = 0 ; idx < nbSubCellGroups ; ++idx){
......@@ -555,7 +557,6 @@ __kernel void FOpenCL__transferInoutPassPerformMpi(__global unsigned char* curr
__global unsigned char* externalCellsPtr, size_t externalCellsSize,
int idxLevel, const __global OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, __global void* userkernel){
return;// TODO
struct FOpenCLGroupOfCells currentCells = BuildFOpenCLGroupOfCells(currentCellsPtr, currentCellsSize);
struct FOpenCLGroupOfCells cellsOther = BuildFOpenCLGroupOfCells(externalCellsPtr, externalCellsSize);
......@@ -585,7 +586,6 @@ __kernel void FOpenCL__transferInoutPassPerformMpi(__global unsigned char* curr
__kernel void FOpenCL__transferInPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
int idxLevel, __global void* userkernel){
return;// TODO
struct FOpenCLGroupOfCells currentCells = BuildFOpenCLGroupOfCells(currentCellsPtr, currentCellsSize);
const MortonIndex blockStartIdx = FOpenCLGroupOfCells_getStartingIndex(&currentCells);
......@@ -627,7 +627,6 @@ __kernel void FOpenCL__transferInoutPassPerform(__global unsigned char* currentC
__global unsigned char* externalCellsPtr, size_t externalCellsSize,
int idxLevel, const __global OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, __global void* userkernel){
return;// TODO
struct FOpenCLGroupOfCells currentCells = BuildFOpenCLGroupOfCells(currentCellsPtr, currentCellsSize);
struct FOpenCLGroupOfCells cellsOther = BuildFOpenCLGroupOfCells(externalCellsPtr, externalCellsSize);
......@@ -662,7 +661,6 @@ __kernel void FOpenCL__transferInoutPassPerform(__global unsigned char* currentC
__kernel void FOpenCL__downardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize,
int nbSubCellGroups, int idxLevel, __global void* userkernel){
return;// TODO
FOpenCLAssertLF(nbSubCellGroups != 0);
struct FOpenCLGroupOfCells currentCells = BuildFOpenCLGroupOfCells(currentCellsPtr, currentCellsSize);
struct FOpenCLGroupOfCells subCellGroups[9];
......@@ -710,7 +708,6 @@ __kernel void FOpenCL__directInoutPassPerformMpi(__global unsigned char* contain
__global unsigned char* externalContainersPtr, size_t externalContainersSize,
const __global OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){
return;// TODO
struct FOpenCLGroupOfParticles containers = BuildFOpenCLGroupOfParticles(containersPtr, containersSize);
struct FOpenCLGroupOfParticles containersOther = BuildFOpenCLGroupOfParticles(externalContainersPtr, externalContainersSize);
......@@ -734,7 +731,6 @@ __kernel void FOpenCL__directInoutPassPerformMpi(__global unsigned char* contain
__kernel void FOpenCL__directInPassPerform(__global unsigned char* containersPtr, size_t containersSize,
const int treeHeight, __global void* userkernel){
return;// TODO
struct FOpenCLGroupOfParticles containers = BuildFOpenCLGroupOfParticles(containersPtr, containersSize);
const MortonIndex blockStartIdx = FOpenCLGroupOfParticles_getStartingIndex(&containers);
......@@ -773,7 +769,6 @@ __kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containers
__global unsigned char* externalContainersPtr, size_t externalContainersSize,
const __global OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){
return;// TODO
struct FOpenCLGroupOfParticles containers = BuildFOpenCLGroupOfParticles(containersPtr, containersSize);
struct FOpenCLGroupOfParticles containersOther = BuildFOpenCLGroupOfParticles(externalContainersPtr, externalContainersSize);
......@@ -801,7 +796,6 @@ __kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containers
__kernel void FOpenCL__mergePassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize,
__global unsigned char* containersPtr, size_t containersSize,
__global void* userkernel){
return;// TODO
struct FOpenCLGroupOfCells leafCells = BuildFOpenCLGroupOfCells(leafCellsPtr,leafCellsSize);
struct FOpenCLGroupOfParticles containers = BuildFOpenCLGroupOfParticles(containersPtr,containersSize);
......
......@@ -160,7 +160,7 @@ int main(int argc, char* argv[]){
// Run the algorithm
GroupKernelClass groupkernel;
GroupAlgorithm groupalgo(&groupedTree,&groupkernel);
groupalgo.execute(FFmmP2M); // TODO
groupalgo.execute(); // FFmmP2M TODO
// Usual algorithm
KernelClass kernels; // FTestKernels FBasicKernels
......
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