Commit 8de14da8 authored by BRAMAS Berenger's avatar BRAMAS Berenger

Make OpenCL compile (but there is still a problem of back copy to the CPU)

parent cb33f0fc
......@@ -5,7 +5,7 @@ typedef long long int MortonIndex;
#define DefaultStructAlign ___DefaultStructAlign___
typedef struct OutOfBlockInteraction{
struct OutOfBlockInteraction{
MortonIndex outIndex;
MortonIndex insideIndex;
int outPosition;
......@@ -16,46 +16,49 @@ struct Uptr9{
struct size_t9{
size_t v[9];
}__attribute__ ((aligned (DefaultStructAlign)));
__kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize,
__kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize,__global unsigned char* leafCellsUpPtr,
__global unsigned char* containersPtr, size_t containersSize,
__global void* userkernel ){
}
__kernel void FOpenCL__upwardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize,
__kernel void FOpenCL__upwardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize, __global unsigned char* currentCellsUpPtr,
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize, struct Uptr9 subCellGroupsUpPtr,
int nbSubCellGroups, int idxLevel, __global void* userkernel){
}
__kernel void FOpenCL__transferInoutPassPerformMpi(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
__global unsigned char* externalCellsPtr, size_t externalCellsSize,
__kernel void FOpenCL__transferInoutPassPerformMpi(__global unsigned char* currentCellsPtr, size_t currentCellsSize, __global unsigned char* currentCellsDownPtr,
__global unsigned char* externalCellsPtr, size_t externalCellsSize, __global unsigned char* externalCellsUpPtr,
int idxLevel, const __global struct OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, __global void* userkernel){
}
__kernel void FOpenCL__transferInPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
__global unsigned char* currentCellsUpPtr, __global unsigned char* currentCellsDownPtr,
int idxLevel, __global void* userkernel){
}
__kernel void FOpenCL__transferInoutPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
__global unsigned char* currentCellsUpPtr, __global unsigned char* currentCellsDownPtr,
__global unsigned char* externalCellsPtr, size_t externalCellsSize,
__global unsigned char* externalCellsUpPtr, __global unsigned char* externalCellsDownPtr,
int idxLevel, const __global struct OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, __global void* userkernel){
}
__kernel void FOpenCL__downardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize,
__kernel void FOpenCL__downardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize, __global unsigned char* currentCellsDownPtr,
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize, struct Uptr9 subCellGroupsDownPtr,
int nbSubCellGroups, int idxLevel, __global void* userkernel){
}
__kernel void FOpenCL__directInoutPassPerformMpi(__global unsigned char* containersPtr, size_t containersSize,
__global unsigned char* externalContainersPtr, size_t externalContainersSize,
__kernel void FOpenCL__directInoutPassPerformMpi(__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,
__global unsigned char* externalContainersPtr, size_t externalContainersSize, __global unsigned char* outsideInteractionsCl,
const __global struct OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){
}
__kernel void FOpenCL__directInPassPerform(__global unsigned char* containersPtr, size_t containersSize,
__kernel void FOpenCL__directInPassPerform(__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,
const int treeHeight, __global void* userkernel){
}
__kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containersPtr, size_t containersSize,
__global unsigned char* externalContainersPtr, size_t externalContainersSize,
__kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,
__global unsigned char* externalContainersPtr, size_t externalContainersSize, __global unsigned char* externalContainersDownPtr,
const __global struct OutOfBlockInteraction* outsideInteractions,
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){
}
__kernel void FOpenCL__mergePassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize,
__global unsigned char* containersPtr, size_t containersSize,
__kernel void FOpenCL__mergePassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize, __global unsigned char* leafCellsDownPtr,
__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,
__global void* userkernel){
}
......@@ -14,61 +14,65 @@ public:
const char* getKernelCode(const int /*inDevId*/){
const char* kernelcode =
"typedef long long int MortonIndex; \
#define DefaultStructAlign " FStarPUDefaultAlignStr "\
typedef struct OutOfBlockInteraction{ \
MortonIndex outIndex; \
MortonIndex insideIndex; \
int outPosition; \
} __attribute__ ((aligned (DefaultStructAlign))); \
struct Uptr9{ \
__global unsigned char* ptrs[9]; \
} __attribute__ ((aligned (DefaultStructAlign))); \
struct size_t9{ \
size_t v[9]; \
}__attribute__ ((aligned (DefaultStructAlign))); \
__kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize, \
__global unsigned char* containersPtr, size_t containersSize, \
__global void* userkernel ){ \
} \
__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){ \
} \
__kernel void FOpenCL__transferInoutPassPerformMpi(__global unsigned char* currentCellsPtr, size_t currentCellsSize, \
__global unsigned char* externalCellsPtr, size_t externalCellsSize, \
int idxLevel, const __global struct OutOfBlockInteraction* outsideInteractions, \
size_t nbOutsideInteractions, __global void* userkernel){ \
} \
__kernel void FOpenCL__transferInPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize, \
int idxLevel, __global void* userkernel){ \
} \
__kernel void FOpenCL__transferInoutPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize, \
__global unsigned char* externalCellsPtr, size_t externalCellsSize, \
int idxLevel, const __global struct OutOfBlockInteraction* outsideInteractions, \
size_t nbOutsideInteractions, __global void* userkernel){ \
} \
__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){ \
} \
__kernel void FOpenCL__directInoutPassPerformMpi(__global unsigned char* containersPtr, size_t containersSize, \
__global unsigned char* externalContainersPtr, size_t externalContainersSize, \
const __global struct OutOfBlockInteraction* outsideInteractions, \
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){ \
} \
__kernel void FOpenCL__directInPassPerform(__global unsigned char* containersPtr, size_t containersSize, \
const int treeHeight, __global void* userkernel){ \
} \
__kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containersPtr, size_t containersSize, \
__global unsigned char* externalContainersPtr, size_t externalContainersSize, \
const __global struct OutOfBlockInteraction* outsideInteractions, \
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){ \
} \
__kernel void FOpenCL__mergePassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize, \
__global unsigned char* containersPtr, size_t containersSize, \
__global void* userkernel){ \
}";
"typedef long long int MortonIndex; \
#define DefaultStructAlign " FStarPUDefaultAlignStr "\
\
typedef struct OutOfBlockInteraction{\
MortonIndex outIndex;\
MortonIndex insideIndex;\
int outPosition;\
} __attribute__ ((aligned (DefaultStructAlign)));\
struct Uptr9{\
__global unsigned char* ptrs[9];\
} __attribute__ ((aligned (DefaultStructAlign)));\
struct size_t9{\
size_t v[9];\
}__attribute__ ((aligned (DefaultStructAlign)));\
__kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize,__global unsigned char* leafCellsUpPtr,\
__global unsigned char* containersPtr, size_t containersSize,\
__global void* userkernel ){\
}\
__kernel void FOpenCL__upwardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize, __global unsigned char* currentCellsUpPtr,\
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize, struct Uptr9 subCellGroupsUpPtr,\
int nbSubCellGroups, int idxLevel, __global void* userkernel){\
}\
__kernel void FOpenCL__transferInoutPassPerformMpi(__global unsigned char* currentCellsPtr, size_t currentCellsSize, __global unsigned char* currentCellsDownPtr,\
__global unsigned char* externalCellsPtr, size_t externalCellsSize, __global unsigned char* externalCellsUpPtr,\
int idxLevel, const __global struct OutOfBlockInteraction* outsideInteractions,\
size_t nbOutsideInteractions, __global void* userkernel){\
}\
__kernel void FOpenCL__transferInPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,\
__global unsigned char* currentCellsUpPtr, __global unsigned char* currentCellsDownPtr,\
int idxLevel, __global void* userkernel){\
}\
__kernel void FOpenCL__transferInoutPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize,\
__global unsigned char* currentCellsUpPtr, __global unsigned char* currentCellsDownPtr,\
__global unsigned char* externalCellsPtr, size_t externalCellsSize,\
__global unsigned char* externalCellsUpPtr, __global unsigned char* externalCellsDownPtr,\
int idxLevel, const __global struct OutOfBlockInteraction* outsideInteractions,\
size_t nbOutsideInteractions, __global void* userkernel){\
}\
__kernel void FOpenCL__downardPassPerform(__global unsigned char* currentCellsPtr, size_t currentCellsSize, __global unsigned char* currentCellsDownPtr,\
struct Uptr9 subCellGroupsPtr, struct size_t9 subCellGroupsSize, struct Uptr9 subCellGroupsDownPtr,\
int nbSubCellGroups, int idxLevel, __global void* userkernel){\
}\
__kernel void FOpenCL__directInoutPassPerformMpi(__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,\
__global unsigned char* externalContainersPtr, size_t externalContainersSize, __global unsigned char* outsideInteractionsCl,\
const __global struct OutOfBlockInteraction* outsideInteractions,\
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){\
}\
__kernel void FOpenCL__directInPassPerform(__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,\
const int treeHeight, __global void* userkernel){\
}\
__kernel void FOpenCL__directInoutPassPerform(__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,\
__global unsigned char* externalContainersPtr, size_t externalContainersSize, __global unsigned char* externalContainersDownPtr,\
const __global struct OutOfBlockInteraction* outsideInteractions,\
size_t nbOutsideInteractions, const int treeHeight, __global void* userkernel){\
}\
__kernel void FOpenCL__mergePassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize, __global unsigned char* leafCellsDownPtr,\
__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersDownPtr,\
__global void* userkernel){\
}";
return kernelcode;
}
......
......@@ -157,7 +157,8 @@ public:
}
void bottomPassPerform(cl_mem leafCellsPtr, size_t leafCellsSize, cl_mem leafCellsUpPtr, cl_mem containersPtr, size_t containersSize){
SetKernelArgs(kernel_bottomPassPerform, 0, &leafCellsPtr, &leafCellsSize, &leafCellsUpPtr, &containersPtr, &containersSize, &user_data/*, &outputcl*/);
SetKernelArgs(kernel_bottomPassPerform, 0, &leafCellsPtr, &leafCellsSize, &leafCellsUpPtr,
&containersPtr, &containersSize, &user_data/*, &outputcl*/);
const int err = clEnqueueNDRangeKernel(queue_bottomPassPerform, kernel_bottomPassPerform, kernelFilename.getNbDims(), NULL,
kernelFilename.getNbGroups(), kernelFilename.getGroupSize(), 0, NULL, NULL);
if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
......
This diff is collapsed.
......@@ -12,10 +12,11 @@ class FTestOpenCLCode{
public:
//FTestOpenCLCode() : kernelfile("/home/berenger/Projets/ScalfmmGit/scalfmm/Src/GroupTree/OpenCl/FEmptyKernel.cl"){
FTestOpenCLCode() : kernelfile("/home/berenger/Projets/ScalfmmGit/scalfmm/Src/GroupTree/OpenCl/FTestKernel.cl"){
FTestOpenCLCode() : kernelfile("/home/berenger/Projets/ScalfmmGit/scalfmm/Src/GroupTree/TestKernel/FTestKernel.cl"){
kernelfile.replaceAll("___FReal___", "double");
kernelfile.replaceAll("___FParticleValueClass___", "long long");
kernelfile.replaceAll("___NbAttributesPerParticle___", 2);
kernelfile.replaceAll("___NbSymbAttributes___", 0);
kernelfile.replaceAll("___NbAttributesPerParticle___", 1);
const size_t structAlign = FStarPUDefaultAlign::StructAlign;
kernelfile.replaceAll("___DefaultStructAlign___", structAlign);
......
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