FEmptyOpenCLCode.hpp 6.34 KB
Newer Older
1 2 3 4 5 6
// @SCALFMM_PRIVATE
#ifndef FEMPTYOPENCLCODE_HPP
#define FEMPTYOPENCLCODE_HPP

// Return the same thing as FEmptyKernel.cl

7
#include "../StarPUUtils/FStarPUDefaultAlign.hpp"
BRAMAS Berenger's avatar
BRAMAS Berenger committed
8

9 10 11 12 13 14 15 16
class FEmptyOpenCLCode{

public:
    FEmptyOpenCLCode(){
    }

    const char* getKernelCode(const int /*inDevId*/){
        const char* kernelcode =
17 18 19
                "typedef long long int MortonIndex; \
                #define DefaultStructAlign " FStarPUDefaultAlignStr "\
                \
20
                struct OutOfBlockInteraction{\
21 22
                    MortonIndex outIndex;\
                    MortonIndex insideIndex;\
BRAMAS Berenger's avatar
BRAMAS Berenger committed
23
                    int relativeOutPosition;\
24
                    int insideIdxInBlock;\
25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76
                } __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){\
                }";
77 78 79 80 81 82
        return kernelcode;
    }

    void releaseKernelCode(){
    }

BRAMAS Berenger's avatar
BRAMAS Berenger committed
83
    unsigned int getNbDims() const {
84 85 86
        return 0;
    }

87
    const size_t* getNbGroups(const int /*inSizeInterval*/) const {
88 89 90 91
        return nullptr;
    }

    const size_t* getGroupSize() const {
92 93 94 95 96 97
        return nullptr;
    }
};

#endif // FEMPTYOPENCLCODE_HPP