MAJ terminée. Nous sommes passés en version 14.6.2 . Pour consulter les "releases notes" associées c'est ici :

https://about.gitlab.com/releases/2022/01/11/security-release-gitlab-14-6-2-released/
https://about.gitlab.com/releases/2022/01/04/gitlab-14-6-1-released/

Commit 444eaafe authored by BRAMAS Berenger's avatar BRAMAS Berenger
Browse files

Wip starpu opencl test kernel

parent 487d31b2
......@@ -117,9 +117,9 @@ public:
// To get access to descriptor
friend struct FTestCellDescriptor;
friend struct FTestCell_Alignement;
};
#endif //FTESTCELL_HPP
/** This file contains the prototype for a kernel in opencl */
// @SCALFMM_PRIVATE
typedef long long int MortonIndex;
......@@ -20,10 +21,6 @@ typedef struct OutOfBlockInteraction_t{
/***************************************************************************/
/***************************************************************************/
struct Uptr8{
__global unsigned char* ptrs[8];
};
struct Uptr9{
__global unsigned char* ptrs[9];
};
......
......@@ -138,7 +138,6 @@ public:
void upwardPassPerform(cl_mem currentCellsPtr, size_t currentCellsSize, cl_mem subCellGroupsPtr[9], size_t subCellGroupsSize[9], int nbSubCellGroups, int idxLevel){
return;
Uptr9 ptrs;
memcpy(ptrs.ptrs, subCellGroupsPtr, sizeof(cl_mem)*9);
size_t9 sizes;
......
/** This file contains the prototype for a kernel in opencl */
// @SCALFMM_PRIVATE
typedef ___FReal___ FReal;
typedef ___FParticleValueClass___ FParticleValueClass;
typedef long long int MortonIndex;
#define FCellClassSize ___FCellClassSize___
#define FCellUpOffset ___FCellUpOffset___
#define FCellDownOffset ___FCellDownOffset___
#define FOpenCLGroupOfCellsCellIsEmptyFlag ((MortonIndex)-1)
#define NbAttributesPerParticle ___NbAttributesPerParticle___
......@@ -377,11 +379,6 @@ __global unsigned char* FOpenCLGroupOfCells_getCell(struct FOpenCLGroupOfCells*
/***************************************************************************/
/***************************************************************************/
/***************************************************************************/
struct Uptr8{
__global unsigned char* ptrs[8];
};
struct Uptr9{
__global unsigned char* ptrs[9];
};
......@@ -395,6 +392,8 @@ struct Uptr343{
};
void P2M(__global unsigned char* pole, const struct FOpenCLGroupAttachedLeaf particles, __global void* user_data) {
__global long long* up = (__global long long*)(pole+FCellUpOffset);
(*up) = particles.nbParticles;
}
void M2M(__global unsigned char* pole, __global unsigned char* child[8], const int level, __global void* user_data) {
......@@ -476,6 +475,7 @@ __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){
......@@ -523,6 +523,7 @@ __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);
......@@ -552,6 +553,7 @@ __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);
......@@ -593,6 +595,7 @@ __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);
......@@ -627,6 +630,7 @@ __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];
......@@ -674,6 +678,7 @@ __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);
......@@ -697,6 +702,7 @@ __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);
......@@ -735,6 +741,7 @@ __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);
......@@ -762,6 +769,7 @@ __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);
......
#ifndef FTEXTREPLACER_HPP
#define FTEXTREPLACER_HPP
// @SCALFMM_PRIVATE
#include "../../Utils/FGlobal.hpp"
#include "../../Utils/FAssert.hpp"
......
......@@ -39,6 +39,15 @@
#include "../../Src/GroupTree/FStarPUKernelCapacities.hpp"
#include "../../Src/GroupTree/OpenCl/FTextReplacer.hpp"
struct FTestCell_Alignement{
static const int dataUp;
static const int dataDown;
};
const int FTestCell_Alignement::dataUp = reinterpret_cast<std::size_t>(&((reinterpret_cast<FTestCell*>(0xF00))->dataUp)) - std::size_t(0xF00);
const int FTestCell_Alignement::dataDown = reinterpret_cast<std::size_t>(&((reinterpret_cast<FTestCell*>(0xF00))->dataDown)) - std::size_t(0xF00);
int main(int argc, char* argv[]){
const FParameterNames LocalOptionBlocSize {
{"-bs"},
......@@ -53,11 +62,14 @@ int main(int argc, char* argv[]){
FTextReplacer kernelfile;
public:
OpenCLSource() : kernelfile("/home/berenger/Projets/ScalfmmGit/scalfmm/Src/GroupTree/OpenCl/FEmptyKernel.cl"){
//OpenCLSource() : kernelfile("/home/berenger/Projets/ScalfmmGit/scalfmm/Src/GroupTree/OpenCl/FEmptyKernel.cl"){
OpenCLSource() : kernelfile("/home/berenger/Projets/ScalfmmGit/scalfmm/Src/GroupTree/OpenCl/FTestKernel.cl"){
kernelfile.replaceAll("___FReal___", "double");
kernelfile.replaceAll("___FParticleValueClass___", "double");
kernelfile.replaceAll("___FCellClassSize___", 4);
kernelfile.replaceAll("___NbAttributesPerParticle___", 4);
kernelfile.replaceAll("___FParticleValueClass___", "long long");
kernelfile.replaceAll("___FCellClassSize___", sizeof(FTestCell));
kernelfile.replaceAll("___NbAttributesPerParticle___", 2);
kernelfile.replaceAll("___FCellUpOffset___", FTestCell_Alignement::dataUp);
kernelfile.replaceAll("___FCellDownOffset___", FTestCell_Alignement::dataDown);
}
operator const char*(){
......
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