FCudaGroupAttachedLeaf.hpp 6.34 KB
Newer Older
1 2 3 4 5
#ifndef FCUDAGROUPATTACHEDLEAF_HPP
#define FCUDAGROUPATTACHEDLEAF_HPP

#include "FCudaGlobal.hpp"

6
template <class FReal, unsigned NbSymbAttributes, unsigned NbAttributesPerParticle, class AttributeClass = FReal>
7 8 9
class FCudaGroupAttachedLeaf {
protected:
    //< Nb of particles in the current leaf
10
    FSize nbParticles;
11 12 13
    //< Pointers to the positions of the particles
    FReal* positionsPointers[3];
    //< Pointers to the attributes of the particles
14
    AttributeClass* attributes[NbSymbAttributes+NbAttributesPerParticle];
15 16 17 18 19

public:
    /** Empty constructor to point to nothing */
    __device__ FCudaGroupAttachedLeaf() : nbParticles(-1) {
        memset(positionsPointers, 0, sizeof(FReal*) * 3);
20
        memset(attributes, 0, sizeof(AttributeClass*) * (NbSymbAttributes+NbAttributesPerParticle));
21 22 23 24 25 26 27 28 29 30
    }

    /**
     * @brief FCudaGroupAttachedLeaf
     * @param inNbParticles the number of particles in the leaf
     * @param inPositionBuffer the memory address of the X array of particls
     * @param inLeadingPosition each position is access by inPositionBuffer + in bytes inLeadingPosition*idx
     * @param inAttributesBuffer the memory address of the first attribute
     * @param inLeadingAttributes each attribute is access by inAttributesBuffer + in bytes inLeadingAttributes*idx
     */
31
    __device__ FCudaGroupAttachedLeaf(const FSize inNbParticles, FReal* inPositionBuffer, const size_t inLeadingPosition,
32 33 34 35 36 37 38
                       AttributeClass* inAttributesBuffer, const size_t inLeadingAttributes)
        : nbParticles(inNbParticles){
        // Redirect pointers to position
        positionsPointers[0] = inPositionBuffer;
        positionsPointers[1] = reinterpret_cast<FReal*>(reinterpret_cast<unsigned char*>(inPositionBuffer) + inLeadingPosition);
        positionsPointers[2] = reinterpret_cast<FReal*>(reinterpret_cast<unsigned char*>(inPositionBuffer) + inLeadingPosition*2);

39
        for(unsigned idxAttribute = 0 ; idxAttribute < NbSymbAttributes ; ++idxAttribute){
40
            attributes[idxAttribute] = reinterpret_cast<AttributeClass*>(reinterpret_cast<unsigned char*>(inPositionBuffer) + inLeadingPosition*3 + inLeadingAttributes*idxAttribute);
41 42
        }

43
        // Redirect pointers to data
44 45 46 47 48 49 50
        if(inAttributesBuffer){
            for(unsigned idxAttribute = 0 ; idxAttribute < NbAttributesPerParticle ; ++idxAttribute){
                attributes[idxAttribute+NbSymbAttributes] = reinterpret_cast<AttributeClass*>(reinterpret_cast<unsigned char*>(inAttributesBuffer) + idxAttribute*inLeadingAttributes);
            }
        }
        else{
            memset(&attributes[NbSymbAttributes], 0, sizeof(AttributeClass*)*NbAttributesPerParticle);
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 77 78 79 80 81 82 83 84 85
        }
    }

    /** Copy the attached group to another one (copy the pointer not the content!) */
    __device__ FCudaGroupAttachedLeaf(const FCudaGroupAttachedLeaf& other) : nbParticles(other.nbParticles) {
        positionsPointers[0] = other.positionsPointers[0];
        positionsPointers[1] = other.positionsPointers[1];
        positionsPointers[2] = other.positionsPointers[2];

        // Redirect pointers to data
        for(unsigned idxAttribute = 0 ; idxAttribute < NbAttributesPerParticle ; ++idxAttribute){
            attributes[idxAttribute] = other.attributes[idxAttribute];
        }
    }

    /** Copy the attached group to another one (copy the pointer not the content!) */
    __device__ FCudaGroupAttachedLeaf& operator=(const FCudaGroupAttachedLeaf& other){
        nbParticles = (other.nbParticles);

        positionsPointers[0] = other.positionsPointers[0];
        positionsPointers[1] = other.positionsPointers[1];
        positionsPointers[2] = other.positionsPointers[2];

        // Redirect pointers to data
        for(unsigned idxAttribute = 0 ; idxAttribute < NbAttributesPerParticle ; ++idxAttribute){
            attributes[idxAttribute] = other.attributes[idxAttribute];
        }

        return (*this);
    }

    /**
     * @brief getNbParticles
     * @return the number of particles in the leaf
     */
86
    __device__ FSize getNbParticles() const{
87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148
        return nbParticles;
    }

    /**
     * @brief getPositions
     * @return a FReal*[3] to get access to the positions
     */
    __device__ const FReal*const* getPositions() const {
        return positionsPointers;
    }

    /**
     * @brief getWPositions
     * @return get the position in write mode
     */
    __device__ FReal* const* getWPositions() {
        return positionsPointers;
    }

    /**
     * @brief getAttribute
     * @param index
     * @return the attribute at index index
     */
    __device__ AttributeClass* getAttribute(const int index) {
        return attributes[index];
    }

    /**
     * @brief getAttribute
     * @param index
     * @return
     */
    __device__ const AttributeClass* getAttribute(const int index) const {
        return attributes[index];
    }

    /**
     * Get the attribute with a forcing compile optimization
     */
    template <int index>
    __device__ AttributeClass* getAttribute() {
        static_assert(index < NbAttributesPerParticle, "Index to get attributes is out of scope.");
        return attributes[index];
    }

    /**
     * Get the attribute with a forcing compile optimization
     */
    template <int index>
    __device__ const AttributeClass* getAttribute() const {
        static_assert(index < NbAttributesPerParticle, "Index to get attributes is out of scope.");
        return attributes[index];
    }

    /** Return true if it has been attached to a memoy block */
    __device__ bool isAttachedToSomething() const {
        return nbParticles != -1;
    }

    /** Copy data for one particle (from the ParticleClassContainer into the attached buffer) */
    template<class ParticleClassContainer>
149
    __device__ void setParticle(const FSize destPartIdx, const FSize srcPartIdx, const ParticleClassContainer* particles){
150 151 152 153 154 155 156 157 158 159 160 161 162 163
        // Copy position
        positionsPointers[0][destPartIdx] = particles->getPositions()[0][srcPartIdx];
        positionsPointers[1][destPartIdx] = particles->getPositions()[1][srcPartIdx];
        positionsPointers[2][destPartIdx] = particles->getPositions()[2][srcPartIdx];

        // Copy data
        for(unsigned idxAttribute = 0 ; idxAttribute < NbAttributesPerParticle ; ++idxAttribute){
            attributes[idxAttribute][destPartIdx] = particles->getAttribute(idxAttribute)[srcPartIdx];
        }
    }
};

#endif // FCUDAGROUPATTACHEDLEAF_HPP