Squishy.Sim is an open source project powered by Assembla

Assembla offers free public and private SVN/Git repositories and project hosting with bug/issue tracking and collaboration tools.

Commit 7838344a88937312f99744af850f8b2fba4190c4

User picture

- first version of z indexing on CPU and can toggle now between GPUand CPU solver

Files Affected

 
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
 
 
1
#ifndef SPH_CLUTIL_H
 
 
2
#define SPH_CLUTIL_H
 
 
3
 
 
 
4
#include <CL/cl.h>
 
 
5
 
 
 
6
namespace Squishy { namespace CL {
 
 
7
    const char* GetCLErrorString(cl_int err) {
 
 
8
        switch (err) {
 
 
9
        case CL_SUCCESS:                            return "Success!";
 
 
10
        case CL_DEVICE_NOT_FOUND:                   return "Device not found";
 
 
11
        case CL_DEVICE_NOT_AVAILABLE:               return "Device not available";
 
 
12
        case CL_COMPILER_NOT_AVAILABLE:             return "Compiler not available";
 
 
13
        case CL_MEM_OBJECT_ALLOCATION_FAILURE:      return "Memory object allocation failure";
 
 
14
        case CL_OUT_OF_RESOURCES:                   return "Out of resources";
 
 
15
        case CL_OUT_OF_HOST_MEMORY:                 return "Out of host memory";
 
 
16
        case CL_PROFILING_INFO_NOT_AVAILABLE:       return "Profiling information not available";
 
 
17
        case CL_MEM_COPY_OVERLAP:                   return "Memory copy overlap";
 
 
18
        case CL_IMAGE_FORMAT_MISMATCH:              return "Image format mismatch";
 
 
19
        case CL_IMAGE_FORMAT_NOT_SUPPORTED:         return "Image format not supported";
 
 
20
        case CL_BUILD_PROGRAM_FAILURE:              return "Program build failure";
 
 
21
        case CL_MAP_FAILURE:                        return "Map failure";
 
 
22
        case CL_INVALID_VALUE:                      return "Invalid value";
 
 
23
        case CL_INVALID_DEVICE_TYPE:                return "Invalid device type";
 
 
24
        case CL_INVALID_PLATFORM:                   return "Invalid platform";
 
 
25
        case CL_INVALID_DEVICE:                     return "Invalid device";
 
 
26
        case CL_INVALID_CONTEXT:                    return "Invalid context";
 
 
27
        case CL_INVALID_QUEUE_PROPERTIES:           return "Invalid queue properties";
 
 
28
        case CL_INVALID_COMMAND_QUEUE:              return "Invalid command queue";
 
 
29
        case CL_INVALID_HOST_PTR:                   return "Invalid host pointer";
 
 
30
        case CL_INVALID_MEM_OBJECT:                 return "Invalid memory object";
 
 
31
        case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:    return "Invalid image format descriptor";
 
 
32
        case CL_INVALID_IMAGE_SIZE:                 return "Invalid image size";
 
 
33
        case CL_INVALID_SAMPLER:                    return "Invalid sampler";
 
 
34
        case CL_INVALID_BINARY:                     return "Invalid binary";
 
 
35
        case CL_INVALID_BUILD_OPTIONS:              return "Invalid build options";
 
 
36
        case CL_INVALID_PROGRAM:                    return "Invalid program";
 
 
37
        case CL_INVALID_PROGRAM_EXECUTABLE:         return "Invalid program executable";
 
 
38
        case CL_INVALID_KERNEL_NAME:                return "Invalid kernel name";
 
 
39
        case CL_INVALID_KERNEL_DEFINITION:          return "Invalid kernel definition";
 
 
40
        case CL_INVALID_KERNEL:                     return "Invalid kernel";
 
 
41
        case CL_INVALID_ARG_INDEX:                  return "Invalid argument index";
 
 
42
        case CL_INVALID_ARG_VALUE:                  return "Invalid argument value";
 
 
43
        case CL_INVALID_ARG_SIZE:                   return "Invalid argument size";
 
 
44
        case CL_INVALID_KERNEL_ARGS:                return "Invalid kernel arguments";
 
 
45
        case CL_INVALID_WORK_DIMENSION:             return "Invalid work dimension";
 
 
46
        case CL_INVALID_WORK_GROUP_SIZE:            return "Invalid work group size";
 
 
47
        case CL_INVALID_WORK_ITEM_SIZE:             return "Invalid work item size";
 
 
48
        case CL_INVALID_GLOBAL_OFFSET:              return "Invalid global offset";
 
 
49
        case CL_INVALID_EVENT_WAIT_LIST:            return "Invalid event wait list";
 
 
50
        case CL_INVALID_EVENT:                      return "Invalid event";
 
 
51
        case CL_INVALID_OPERATION:                  return "Invalid operation";
 
 
52
        case CL_INVALID_GL_OBJECT:                  return "Invalid OpenGL object";
 
 
53
        case CL_INVALID_BUFFER_SIZE:                return "Invalid buffer size";
 
 
54
        case CL_INVALID_MIP_LEVEL:                  return "Invalid mip-map level";
 
 
55
        default: return "Unknown error";
 
 
56
        }
 
 
57
    }
 
 
58
 
 
 
59
} }
 
 
60
 
 
 
61
#define CL_SUCCESS_OR_ELSE(err, whatElse) \
 
 
62
    if ((err) != CL_SUCCESS) \
 
 
63
    { \
 
 
64
        printf("CL ERROR: %s\n", Squishy::CL::GetCLErrorString(err)); \
 
 
65
        whatElse; \
 
 
66
    }
 
 
67
 
 
 
68
#endif
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
 
 
1
#ifndef SPH_HASHTABLE_H
 
 
2
#define SPH_HASHTABLE_H
 
 
3
 
 
 
4
#include "_SPH_Base.h"
 
 
5
 
 
 
6
#include "DataStructures/SpatialMapping1D.h"
 
 
7
 
 
 
8
namespace Squishy { namespace SPH { namespace DataStructures {
 
 
9
 
 
 
10
    // Reference: "Optimized Spatial Hashing for Collision Detection of Deformable Objects"
 
 
11
    // TODO: Timestamp?
 
 
12
    template<class T>
 
 
13
    class Hashtable3D : public SpatialMapping1D<T>
 
 
14
    {
 
 
15
 
 
 
16
    public:
 
 
17
        Hashtable3D(Real cellLength, int size = 10) : SpatialMapping1D(cellLength, size)
 
 
18
        {
 
 
19
        }
 
 
20
 
 
 
21
        virtual ~Hashtable3D()
 
 
22
        {
 
 
23
        }
 
 
24
 
 
 
25
        virtual int MapIndex(const Vec3i& idx) const
 
 
26
        {
 
 
27
            return (unsigned int)(((idx.x * 73856093) ^ (idx.y * 19349663) ^ (idx.z * 83492791))) % bucketCount;
 
 
28
        }
 
 
29
    };
 
 
30
 
 
 
31
#define ZOrderMaxExtensionBits 10
 
 
32
#define ZOrderMaxExtension 1024
 
 
33
#define ZOrderDimensions 3
 
 
34
typedef int ZOrderValueType;
 
 
35
 
 
 
36
    class ZOrderer // : public SpatialMapping1D<T>
 
 
37
    {
 
 
38
    private:
 
 
39
        ZOrderValueType ZIndexMap[ZOrderMaxExtension][ZOrderDimensions];
 
 
40
        Real cellLength;
 
 
41
        int size;
 
 
42
 
 
 
43
        /**
 
 
44
         * 
 
 
45
         */
 
 
46
        void CreateZIndexMap()
 
 
47
        {
 
 
48
            for (int i = 0; i < ZOrderMaxExtension; ++i)
 
 
49
            {
 
 
50
                for (int k = 0; k < ZOrderDimensions; ++k)
 
 
51
                {
 
 
52
                    int dstIndex = k;
 
 
53
                    for (int j = 0; j < ZOrderMaxExtensionBits; ++j)
 
 
54
                    {
 
 
55
                        //int srcIndex = j;
 
 
56
                        ZOrderValueType mask = 1 << j;
 
 
57
                        int srcBit = (i & mask) != 0;               // compute the value of the j'th bit
 
 
58
                        ZIndexMap[i][k] |= srcBit << dstIndex;      // turn on/off the bit at position: k + j * ZOrderDimensions
 
 
59
 
 
 
60
                        dstIndex += ZOrderDimensions;
 
 
61
                    }
 
 
62
                }
 
 
63
            }
 
 
64
        }
 
 
65
 
 
 
66
    public:
 
 
67
        ZOrderer(Real cellLength, int size = 10) : cellLength(cellLength), size(size)
 
 
68
        {
 
 
69
            CreateZIndexMap();
 
 
70
        }
 
 
71
 
 
 
72
        int GetZIndex(const Vec3i& idx) const
 
 
73
        {
 
 
74
            return ZIndexMap[idx[0]][0] | ZIndexMap[idx[1]][1] | ZIndexMap[idx[2]][2];
 
 
75
        }
 
 
76
 
 
 
77
        int GetZIndex(const Vector& pos) const
 
 
78
        {
 
 
79
            Vec3i idx = Vec3i(glm::floor(pos/cellLength));
 
 
80
            return GetZIndex(idx);
 
 
81
        }
 
 
82
 
 
 
83
        const ZOrderValueType* GetZIndexMap() const { return &ZIndexMap[0][0]; }
 
 
84
    };
 
 
85
} } }
 
 
86
 
 
 
87
#endif
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
 
 
1
#ifndef SPH_SPATIALMAPPING1D_H
 
 
2
#define SPH_SPATIALMAPPING1D_H
 
 
3
 
 
 
4
#include "_SPH_Base.h"
 
 
5
 
 
 
6
#include <vector>
 
 
7
#include <algorithm>
 
 
8
 
 
 
9
namespace Squishy { namespace SPH { namespace DataStructures {
 
 
10
 
 
 
11
        /**
 
 
12
        * Points from the center of an axis-aligned box of length 2 toward its max point.
 
 
13
        */
 
 
14
#define ToMax Vector(1, 1, 1)
 
 
15
 
 
 
16
    // Reference: "Optimized Spatial Hashing for Collision Detection of Deformable Objects"
 
 
17
    // TODO: Timestamp
 
 
18
    template<class T>
 
 
19
    class SpatialMapping1D
 
 
20
    {
 
 
21
    public:
 
 
22
        typedef std::vector<T> Bucket; 
 
 
23
        typedef typename Bucket::iterator BucketIterator;
 
 
24
        typedef typename Bucket::const_iterator BucketIteratorConst;
 
 
25
 
 
 
26
        typedef std::vector<Bucket*> BucketList;
 
 
27
        typedef typename BucketList::iterator BucketListIterator;
 
 
28
 
 
 
29
    protected:
 
 
30
        Real cellLength;
 
 
31
        int bucketCount;
 
 
32
        Bucket* buckets;
 
 
33
 
 
 
34
    public:
 
 
35
        SpatialMapping1D(Real cellLength, int size = 10) : cellLength(cellLength), bucketCount(size)
 
 
36
        {
 
 
37
            buckets = new std::vector<T>[size];
 
 
38
        }
 
 
39
 
 
 
40
        virtual ~SpatialMapping1D()
 
 
41
        {
 
 
42
            delete [] buckets;
 
 
43
        }
 
 
44
 
 
 
45
        inline Real GetCellLength() const { return cellLength; }
 
 
46
        void SetCellLength(Real l) { cellLength = l; }
 
 
47
 
 
 
48
        Vec3i GetIndex(const Vector& pos) const
 
 
49
        {
 
 
50
            // negative floating point numbers are actually rounded up -> Make sure, we always round down:
 
 
51
            return Vec3i(glm::floor(pos / cellLength));
 
 
52
        }
 
 
53
 
 
 
54
        Vector GetPositionMin(const Vec3i& idx) const
 
 
55
        {
 
 
56
            return Vector(idx) * cellLength;
 
 
57
        }
 
 
58
 
 
 
59
        virtual int MapIndex(const Vec3i& idx) const = 0;
 
 
60
 
 
 
61
        int MapPos(const Vector& pos) const
 
 
62
        {
 
 
63
            Vec3i idx = Vec3i(glm::floor(pos/cellLength));
 
 
64
            return MapIndex(idx);
 
 
65
        }
 
 
66
 
 
 
67
        inline Vector GetPosition(T obj) const
 
 
68
        {
 
 
69
            return obj->GetPosition();
 
 
70
        }
 
 
71
 
 
 
72
        inline int GetMap3To1D(T obj) const
 
 
73
        {
 
 
74
            return obj->hash;
 
 
75
        }
 
 
76
 
 
 
77
        inline void SetMap3To1D(T obj, int hash) const
 
 
78
        {
 
 
79
            obj->hash = hash;
 
 
80
        }
 
 
81
 
 
 
82
        inline BucketIterator GetBucketIndex(Bucket& bucket, T obj) const
 
 
83
        {
 
 
84
            //return obj->hashIndex;
 
 
85
            for (BucketIterator it = bucket.begin(); it != bucket.end(); ++it)
 
 
86
            {
 
 
87
                if ((*it)->index == obj->index)
 
 
88
                {
 
 
89
                    return it;
 
 
90
                }
 
 
91
            }
 
 
92
            assert(0);
 
 
93
            return BucketIterator();
 
 
94
        }
 
 
95
 
 
 
96
        /*inline void SetIndex(T obj, int index) const
 
 
97
        {
 
 
98
        obj->hashIndex = index;
 
 
99
        }*/
 
 
100
 
 
 
101
        /** 
 
 
102
        * Adds a point-sized object into the bucket at the given position
 
 
103
        */
 
 
104
        void AddPoint(T obj)
 
 
105
        {
 
 
106
            Vector pos = GetPosition(obj);
 
 
107
            int hash = MapPos(pos);
 
 
108
            AddPoint(obj, hash);
 
 
109
        }
 
 
110
 
 
 
111
        /** 
 
 
112
        * Adds a point-sized object into the bucket at the given position
 
 
113
        */
 
 
114
        void AddPoint(T obj, int hash)
 
 
115
        {
 
 
116
            Bucket& bucket = buckets[hash];
 
 
117
            SetMap3To1D(obj, hash);
 
 
118
            bucket.push_back(obj);
 
 
119
        }
 
 
120
 
 
 
121
        void RemovePoint(T obj, int hash)
 
 
122
        {
 
 
123
            Bucket& bucket = buckets[hash];
 
 
124
            bucket.erase(GetBucketIndex(bucket, obj));
 
 
125
        }
 
 
126
 
 
 
127
        ///** 
 
 
128
        // * Adds a sphere-shaped object with given radius into all buckets around the object's position
 
 
129
        // */
 
 
130
        //void AddSphere(T obj, Real radius)
 
 
131
        //{
 
 
132
        //    Vector pos = GetPosition(obj);
 
 
133
 
 
 
134
        //    Vector minv = pos - radius * ToMax;
 
 
135
        //    Vector maxv = pos + radius * ToMax;
 
 
136
        //    Vec3i minIndex = GetIndex(Vector(minv));
 
 
137
        //    Vec3i maxIndex = GetIndex(Vector(maxv));
 
 
138
 
 
 
139
        //    Vector cellSize(cellLength);
 
 
140
 
 
 
141
        //    Vec3i idx;
 
 
142
        //    Vector current;
 
 
143
        //    for (idx.x = minIndex.x; idx.x <= maxIndex.x; ++idx.x)
 
 
144
        //    {
 
 
145
        //        current.x = idx.x * cellLength;
 
 
146
        //        for (idx.y = minIndex.y; idx.y <= maxIndex.y; ++idx.y)
 
 
147
        //        {
 
 
148
        //            current.y = idx.y * cellLength;
 
 
149
        //            for (idx.z = minIndex.z; idx.z <= maxIndex.z; ++idx.z)
 
 
150
        //            {
 
 
151
        //                current.z = idx.z * cellLength;
 
 
152
 
 
 
153
        //                // TODO: Verify that intersection test is correct
 
 
154
        //                if (IntersectSphereAABB(current, current + cellSize, pos, radius))
 
 
155
        //                {
 
 
156
        //                    Bucket& bucket = buckets[Map3To1D(idx)];
 
 
157
        //                    // TODO: Don't add objects twice to the same bucket
 
 
158
        //                    bucket.push_back(obj);
 
 
159
        //                }
 
 
160
        //            }
 
 
161
        //        }
 
 
162
        //    }
 
 
163
        //}
 
 
164
 
 
 
165
        const Bucket& GetBucket(const Vec3i& idx) const
 
 
166
        {
 
 
167
            return buckets[MapIndex(idx)];
 
 
168
        }
 
 
169
 
 
 
170
        const Bucket& GetBucket(const Vector& pos) const
 
 
171
        {
 
 
172
            return buckets[MapPos(pos)];
 
 
173
        }
 
 
174
 
 
 
175
        /**
 
 
176
        * Gets all buckets that are intersected by a sphere of given radius, centered at the given position.
 
 
177
        */
 
 
178
        void GetBuckets(const Vector& pos, Real radius, BucketList& list) const
 
 
179
        {
 
 
180
            Vector min = pos - radius * ToMax;
 
 
181
            Vector max = pos + radius * ToMax;
 
 
182
 
 
 
183
            Vector current;
 
 
184
            Vector cellSize(cellLength);
 
 
185
 
 
 
186
            int l = (max.x - min.x) / cellLength;
 
 
187
            for (int i = 0; i < l; ++i)
 
 
188
            {
 
 
189
                current.x = min.x + cellLength * i;
 
 
190
                for (int j = 0; j < l; ++j)
 
 
191
                {
 
 
192
                    current.y = min.y + cellLength * j;
 
 
193
                    for (int k = 0; k < l; ++k)
 
 
194
                    {
 
 
195
                        current.z = min.z + cellLength * k;
 
 
196
 
 
 
197
                        if (IntersectSphereAABB(current, current + cellSize, pos, radius))
 
 
198
                        {
 
 
199
                            list.push_back(buckets[Map3To1D(current)]);
 
 
200
                        }
 
 
201
                    }
 
 
202
                }
 
 
203
            }
 
 
204
        }
 
 
205
 
 
 
206
        /**
 
 
207
        * Wether the cell with the given index contains any objects
 
 
208
        */
 
 
209
        bool IsOccupied(const Vec3i& idx) const
 
 
210
        {
 
 
211
            const Bucket& bucket = GetBucket(idx);
 
 
212
 
 
 
213
            //return !bucket.empty();
 
 
214
            // iterate over all particles in the bucket and see whether any of them are inside the box at the given index
 
 
215
            for (BucketIteratorConst it = bucket.begin(); it != bucket.end(); ++it)
 
 
216
            {
 
 
217
                if (IsInCell(*it, idx))
 
 
218
                {
 
 
219
                    return true;
 
 
220
                }
 
 
221
            }
 
 
222
            return false;
 
 
223
        }
 
 
224
 
 
 
225
        /**
 
 
226
        * 
 
 
227
        */
 
 
228
        bool IsInCell(T obj, const Vec3i& idx) const
 
 
229
        {
 
 
230
            Vector minv = GetPositionMin(idx);
 
 
231
            Vector pos = GetPosition(obj) - minv;
 
 
232
 
 
 
233
            return pos.x >= 0 && pos.x <= cellLength &&
 
 
234
                pos.y >= 0 && pos.y <= cellLength &&
 
 
235
                pos.z >= 0 && pos.z <= cellLength;
 
 
236
        }
 
 
237
 
 
 
238
        /**
 
 
239
        * Clears all buckets
 
 
240
        */
 
 
241
        void Clear()
 
 
242
        {
 
 
243
            int i = 0;
 
 
244
            for (Bucket* b = buckets; i < bucketCount; ++i, ++b)
 
 
245
            {
 
 
246
                b->clear();
 
 
247
            }
 
 
248
        }
 
 
249
 
 
 
250
        /**
 
 
251
        * Deletes all buckets and creates a new set of them
 
 
252
        */
 
 
253
        void Reset(int newSize = 10)
 
 
254
        {
 
 
255
            delete [] buckets;
 
 
256
            buckets = new std::vector<T>[bucketCount = newSize];
 
 
257
        }
 
 
258
    };
 
 
259
 
 
 
260
} } }
 
 
261
 
 
 
262
#endif
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
60
 
60
 
61
            ParticleFlags GetFlags() const { return flags; }
61
            ParticleFlags GetFlags() const { return flags; }
62
 
62
 
63
            bool Moves() const { return flags & ParticleFlagsMoveable; }
63
            bool Moves() const { return (flags & ParticleFlagsMoveable) != 0; }
64
            bool IsEnabled() const { return flags & ParticleFlagsEnabled; }
64
            bool IsEnabled() const { return (flags & ParticleFlagsEnabled) != 0; }
65
 
65
 
66
            void SetDensityError(Real value) 
66
            void SetDensityError(Real value) 
67
            { 
67
            { 
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
4
#include "SPHUtil.h"
4
#include "SPHUtil.h"
5
#include "Particle.h"
5
#include "Particle.h"
6
#include "ParticleManipulator.h"
6
#include "ParticleManipulator.h"
7
#include "Hashtable.h"
7
#include "DataStructures/Hashtable.h"
8
 
8
 
9
#include <vector>
9
#include <vector>
10
 
10
 
...
 
...
 
36
        volatile long solving;
36
        volatile long solving;
37
 
37
 
38
        Real kernelRadius, kernelRadiusSq;
38
        Real kernelRadius, kernelRadiusSq;
39
        Squishy::SPH::Hashtable3D<Particle*> hashtable;
39
        Squishy::SPH::DataStructures::Hashtable3D<Particle*> hashtable;
40
        std::vector<std::vector<Particle*> > neighborSet;
40
        std::vector<std::vector<Particle*> > neighborSet;
41
 
41
 
42
        void OnParticlesChanged();
42
        void OnParticlesChanged();
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
4
#include "Solver.h"
4
#include "Solver.h"
5
#include <vector>
5
#include <vector>
6
#include "Kernels.h"
6
#include "Kernels.h"
 
 
7
#include "CLUtil.h"
7
 
8
 
8
typedef struct Materials
9
struct MaterialProperties
9
{
10
{
10
    Real* masses;
11
    Real* masses;
11
    Real* gasStiffness;
12
    Real* gasStiffness;
12
 
13
};
13
} Materials;
 
 
14
 
14
 
15
namespace Squishy { namespace SPH { namespace Solvers {
15
namespace Squishy { namespace SPH { namespace Solvers {
16
 
16
 
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
5
#ifndef SPH_SYSCALLS_H
5
#ifndef SPH_SYSCALLS_H
6
#define SPH_SYSCALLS_H
6
#define SPH_SYSCALLS_H
7
 
7
 
 
 
8
#include <cstdio>
 
 
9
 
8
#ifdef _WIN32
10
#ifdef _WIN32
9
#include <Windows.h>
11
#include <Windows.h>
10
#else
12
#else
...
 
...
 
35
*/
37
*/
36
double GetPerformanceTimeSeconds();
38
double GetPerformanceTimeSeconds();
37
 
39
 
 
 
40
long ComputeFileSize(FILE* f);
 
 
41
 
38
} }
42
} }
39
 
43
 
40
#endif
44
#endif
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
10
#include "Particle.h"
10
#include "Particle.h"
11
#include "ParticleManipulator.h"
11
#include "ParticleManipulator.h"
12
#include "SPHUtil.h"
12
#include "SPHUtil.h"
13
#include "Hashtable.h"
13
#include "DataStructures/Hashtable.h"
14
 
14
 
15
#include "Solvers/Solver.h"
15
#include "Solvers/Solver.h"
16
 
16
 
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
 
 
1
#ifndef SPH_CLOMMON_H
 
 
2
#define SPH_CLOMMON_H
 
 
3
 
 
 
4
 
 
 
5
 
 
 
6
#endif
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
50
    <ClCompile>
50
    <ClCompile>
51
      <WarningLevel>Level4</WarningLevel>
51
      <WarningLevel>Level4</WarningLevel>
52
      <Optimization>Disabled</Optimization>
52
      <Optimization>Disabled</Optimization>
53
      <DisableSpecificWarnings>4800;4819;4512;4251;4201;%(DisableSpecificWarnings)</DisableSpecificWarnings>
53
      <DisableSpecificWarnings>4800;4819;4996;4512;4251;4201;%(DisableSpecificWarnings)</DisableSpecificWarnings>
54
      <PreprocessorDefinitions>_BUILD_SIM_SPH;_WINDLL;_DEBUG;%(PreprocessorDefinitions)</PreprocessorDefinitions>
54
      <PreprocessorDefinitions>_BUILD_SIM_SPH;_WINDLL;_DEBUG;%(PreprocessorDefinitions)</PreprocessorDefinitions>
55
      <AdditionalIncludeDirectories>"$(AMDAPPSDKROOT)\include";</AdditionalIncludeDirectories>
55
      <AdditionalIncludeDirectories>"$(AMDAPPSDKROOT)\include";</AdditionalIncludeDirectories>
56
      <TreatWarningAsError>false</TreatWarningAsError>
56
      <TreatWarningAsError>false</TreatWarningAsError>
...
 
...
 
80
  </ItemDefinitionGroup>
80
  </ItemDefinitionGroup>
81
  <ItemGroup>
81
  <ItemGroup>
82
    <ClCompile Include="src\CollisionHandler.cpp" />
82
    <ClCompile Include="src\CollisionHandler.cpp" />
 
 
83
    <ClCompile Include="src\DataStructures\Hashtable.cpp" />
83
    <ClCompile Include="src\Fluid.cpp" />
84
    <ClCompile Include="src\Fluid.cpp" />
84
    <ClCompile Include="src\Forces.cpp" />
85
    <ClCompile Include="src\Forces.cpp" />
85
    <ClCompile Include="src\Geometry\geomutil.cpp" />
86
    <ClCompile Include="src\Geometry\geomutil.cpp" />
86
    <ClCompile Include="src\Hashtable.cpp" />
 
 
87
    <ClCompile Include="src\Material.cpp" />
87
    <ClCompile Include="src\Material.cpp" />
88
    <ClCompile Include="src\Materials.cpp" />
88
    <ClCompile Include="src\Materials.cpp" />
89
    <ClCompile Include="src\Particle.cpp" />
89
    <ClCompile Include="src\Particle.cpp" />
...
 
...
 
105
    <ClCompile Include="src\util.cpp" />
105
    <ClCompile Include="src\util.cpp" />
106
  </ItemGroup>
106
  </ItemGroup>
107
  <ItemGroup>
107
  <ItemGroup>
 
 
108
    <ClInclude Include="include\CLUtil.h" />
108
    <ClInclude Include="include\CollisionHandler.h" />
109
    <ClInclude Include="include\CollisionHandler.h" />
 
 
110
    <ClInclude Include="include\DataStructures\Hashtable.h" />
 
 
111
    <ClInclude Include="include\DataStructures\SpatialMapping1D.h" />
109
    <ClInclude Include="include\Forces.h" />
112
    <ClInclude Include="include\Forces.h" />
110
    <ClInclude Include="include\Geometry\geomutil.h" />
113
    <ClInclude Include="include\Geometry\geomutil.h" />
111
    <ClInclude Include="include\Geometry\Shapes.h" />
114
    <ClInclude Include="include\Geometry\Shapes.h" />
112
    <ClInclude Include="include\Hashtable.h" />
 
 
113
    <ClInclude Include="include\Materials.h" />
115
    <ClInclude Include="include\Materials.h" />
114
    <ClInclude Include="include\ParticleManipulator.h" />
116
    <ClInclude Include="include\ParticleManipulator.h" />
115
    <ClInclude Include="include\Solvers\PCISPHSolver.h" />
117
    <ClInclude Include="include\Solvers\PCISPHSolver.h" />
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
26
    <Filter Include="OpenCL">
26
    <Filter Include="OpenCL">
27
      <UniqueIdentifier>{705fc1c1-5987-4b0d-bc81-a170a5522497}</UniqueIdentifier>
27
      <UniqueIdentifier>{705fc1c1-5987-4b0d-bc81-a170a5522497}</UniqueIdentifier>
28
    </Filter>
28
    </Filter>
 
 
29
    <Filter Include="include\DataStructures">
 
 
30
      <UniqueIdentifier>{f4917ada-5592-4316-9dde-7bedc14d66ba}</UniqueIdentifier>
 
 
31
    </Filter>
 
 
32
    <Filter Include="src\DataStructures">
 
 
33
      <UniqueIdentifier>{a6479423-f452-4a2c-9394-bad7761872d0}</UniqueIdentifier>
 
 
34
    </Filter>
29
  </ItemGroup>
35
  </ItemGroup>
30
  <ItemGroup>
36
  <ItemGroup>
31
    <ClCompile Include="src\Fluid.cpp">
37
    <ClCompile Include="src\Fluid.cpp">
...
 
...
 
91
    <ClCompile Include="src\SysCalls.cpp">
97
    <ClCompile Include="src\SysCalls.cpp">
92
      <Filter>src</Filter>
98
      <Filter>src</Filter>
93
    </ClCompile>
99
    </ClCompile>
94
    <ClCompile Include="src\Hashtable.cpp">
 
 
95
      <Filter>src</Filter>
 
 
96
    </ClCompile>
 
 
97
    <ClCompile Include="src\Geometry\geomutil.cpp">
100
    <ClCompile Include="src\Geometry\geomutil.cpp">
98
      <Filter>src\Geometry</Filter>
101
      <Filter>src\Geometry</Filter>
99
    </ClCompile>
102
    </ClCompile>
100
    <ClCompile Include="src\Solvers\WSPHGPUSolver.cpp">
103
    <ClCompile Include="src\Solvers\WSPHGPUSolver.cpp">
101
      <Filter>src\Solvers</Filter>
104
      <Filter>src\Solvers</Filter>
102
    </ClCompile>
105
    </ClCompile>
 
 
106
    <ClCompile Include="src\DataStructures\Hashtable.cpp">
 
 
107
      <Filter>src\DataStructures</Filter>
 
 
108
    </ClCompile>
103
  </ItemGroup>
109
  </ItemGroup>
104
  <ItemGroup>
110
  <ItemGroup>
105
    <ClInclude Include="include\Fluid.h">
111
    <ClInclude Include="include\Fluid.h">
...
 
...
 
174
    <ClInclude Include="include\SysCalls.h">
180
    <ClInclude Include="include\SysCalls.h">
175
      <Filter>include</Filter>
181
      <Filter>include</Filter>
176
    </ClInclude>
182
    </ClInclude>
177
    <ClInclude Include="include\Hashtable.h">
 
 
178
      <Filter>include</Filter>
 
 
179
    </ClInclude>
 
 
180
    <ClInclude Include="include\Geometry\geomutil.h">
183
    <ClInclude Include="include\Geometry\geomutil.h">
181
      <Filter>include\Geometry</Filter>
184
      <Filter>include\Geometry</Filter>
182
    </ClInclude>
185
    </ClInclude>
...
 
...
 
186
    <ClInclude Include="include\Geometry\Shapes.h">
189
    <ClInclude Include="include\Geometry\Shapes.h">
187
      <Filter>include\Geometry</Filter>
190
      <Filter>include\Geometry</Filter>
188
    </ClInclude>
191
    </ClInclude>
 
 
192
    <ClInclude Include="include\DataStructures\SpatialMapping1D.h">
 
 
193
      <Filter>include\DataStructures</Filter>
 
 
194
    </ClInclude>
 
 
195
    <ClInclude Include="include\DataStructures\Hashtable.h">
 
 
196
      <Filter>include\DataStructures</Filter>
 
 
197
    </ClInclude>
 
 
198
    <ClInclude Include="include\CLUtil.h">
 
 
199
      <Filter>include</Filter>
 
 
200
    </ClInclude>
189
  </ItemGroup>
201
  </ItemGroup>
190
  <ItemGroup>
202
  <ItemGroup>
191
    <None Include="OpenCL\kernel.cl">
203
    <None Include="OpenCL\kernel.cl">
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
3
#include "System.h"
3
#include "System.h"
4
 
4
 
5
using namespace std;
5
using namespace std;
 
 
6
using namespace Squishy::SPH::DataStructures;
6
 
7
 
7
namespace Squishy { namespace SPH { namespace Solvers {
8
namespace Squishy { namespace SPH { namespace Solvers {
8
 
9
 
...
 
...
 
118
        {
119
        {
119
            Particle* p = *it;
120
            Particle* p = *it;
120
 
121
 
121
            int hash = hashtable.Hash(p->position);
122
            int hash = hashtable.MapPos(p->position);
122
            if (p->hash != hash)
123
            if (p->hash != hash)
123
            {
124
            {
124
                // only update particle cell if it changed
125
                // only update particle cell if it changed
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
1
#include "Solvers/WSPHGPUSolver.h"
1
#include "Solvers/WSPHGPUSolver.h"
2
 
2
 
3
#include "System.h"
3
#include "System.h"
4
#include "CL/cl.h"
4
#include "SysCalls.h"
 
 
5
#include <iostream>
 
 
6
 
 
 
7
#include <CL/cl.h>
5
 
8
 
6
#define MAX_SOURCE_SIZE (0x100000)
 
 
7
#define MAT0 0
9
#define MAT0 0
8
 
10
 
9
using namespace std;
11
using namespace std;
 
 
12
using namespace Squishy::CL;
10
 
13
 
11
namespace Squishy { namespace SPH { namespace Solvers {
14
namespace Squishy { namespace SPH { namespace Solvers {
12
 
15
 
...
 
...
 
39
    {
42
    {
40
        kernelRadius = max(kernelRadius, fluid->GetMaterial()->ComputeKernelRadius(GetConfig().AverageKernelParticleCount));
43
        kernelRadius = max(kernelRadius, fluid->GetMaterial()->ComputeKernelRadius(GetConfig().AverageKernelParticleCount));
41
        kernelRadiusSq = kernelRadius * kernelRadius;
44
        kernelRadiusSq = kernelRadius * kernelRadius;
42
        
45
 
43
        // fix up kernels:
46
        // fix up kernels:
44
        gaussianKernel->SetRadius(kernelRadius);
47
        gaussianKernel->SetRadius(kernelRadius);
45
        viscosityKernel->SetRadius(kernelRadius);
48
        viscosityKernel->SetRadius(kernelRadius);
...
 
...
 
59
 
62
 
60
            particles.push_back(&*p);
63
            particles.push_back(&*p);
61
        }
64
        }
62
        
65
 
63
        ParticlePtrIterator from = particles.begin();
66
        ParticlePtrIterator from = particles.begin();
64
        ParticlePtrIterator to = particles.end();
67
        ParticlePtrIterator to = particles.end();
65
        
68
 
66
        FindNeighbors(from, to);
69
        FindNeighbors(from, to);
67
 
70
 
68
        ComputeDensityPressure(from, to);
71
        ComputeDensityPressure(from, to);
...
 
...
 
85
        //          1. All particle density  2. all particle pressure
88
        //          1. All particle density  2. all particle pressure
86
        //      After invoking the kernel, set the density and pressure of all particles accordingly
89
        //      After invoking the kernel, set the density and pressure of all particles accordingly
87
 
90
 
88
    
91
 
89
        // Load the kernel source code into the array source_str
92
        // Load the kernel source code into the array source_str
90
        FILE *fp;
93
        FILE *fp;
91
        char *source_str;
94
        char *source_str;
92
        size_t source_size;
95
 
93
 
96
        const char kernelFile[] = "../Squishy.SPH/OpenCL/kernel.cl";
94
        fp = fopen("OpenCL/kernel.cl", "r");
97
        fp = fopen(kernelFile, "r");
95
        if (!fp) {
98
        if (!fp)
96
            fprintf(stderr, "Failed to load kernel.\n");
 
 
97
            exit(1);
 
 
98
        }
 
 
99
        source_str = (char*)malloc(MAX_SOURCE_SIZE);
 
 
100
        source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp);
 
 
101
        fclose( fp );
 
 
102
 
 
 
103
        // Get platform and device information
 
 
104
        cl_platform_id platform_id = NULL;
 
 
105
        cl_device_id device_id = NULL;   
 
 
106
        cl_uint ret_num_devices;
 
 
107
        cl_uint ret_num_platforms;
 
 
108
        cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
 
 
109
        ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
 
 
110
 
 
 
111
        // Create an OpenCL context
 
 
112
        cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret);
 
 
113
 
 
 
114
        // Create memory buffers on the device for each vector 
 
 
115
        ParticleVector& particles = system->GetParticles();
 
 
116
        int nParticles = particles.size();
 
 
117
 
 
 
118
        const FluidVector& fluids = system->GetFluids();
 
 
119
        int nFluids = fluids.size();
 
 
120
        Materials materialParam;
 
 
121
        materialParam.masses = new Real[nFluids];
 
 
122
        materialParam.gasStiffness=  new Real[nFluids];
 
 
123
 
 
 
124
        for (int fluid = 0; fluid<nFluids; ++fluid)
 
 
125
        {
99
        {
126
            materialParam.masses[fluid] = fluids[fluid]->GetMaterial()->GetMass();
100
            fprintf(stderr, "Could not open kernel file: %s\n", kernelFile);
127
            materialParam.gasStiffness[fluid] = fluids[fluid]->GetMaterial()->GetGasStiffness();
101
            return;
128
         }
102
        }
 
 
103
 
 
 
104
        // Read kernel into memory
 
 
105
        size_t source_size = ComputeFileSize(fp) + 1;
 
 
106
        source_str = new char[source_size];
 
 
107
        source_size = fread( source_str, 1, source_size, fp);
 
 
108
        fclose( fp );
 
 
109
 
 
 
110
        source_str[source_size] = 0;
 
 
111
 
 
 
112
        // Get platform and device information
 
 
113
        cl_platform_id platform_id = NULL;
 
 
114
        cl_device_id device_id = NULL;   
 
 
115
        cl_uint ret_num_devices;
 
 
116
        cl_uint ret_num_platforms;
 
 
117
 
 
 
118
        cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
 
 
119
        ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &ret_num_devices);
 
 
120
        assert(ret_num_devices > 0);
 
 
121
        CL_SUCCESS_OR_ELSE(ret, return);   
 
 
122
 
 
 
123
        char name[100];
 
 
124
        ret = clGetDeviceInfo (device_id, CL_DEVICE_NAME, sizeof(name), name, NULL);
 
 
125
        CL_SUCCESS_OR_ELSE(ret, return);
 
 
126
        printf ("Device found: %s\n", name);
 
 
127
 
 
 
128
        // Create an OpenCL context
 
 
129
        cl_context_properties contextProperties[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, 0 };
 
 
130
        cl_context context = clCreateContext( contextProperties, 1, &device_id, NULL, NULL, &ret);
 
 
131
        CL_SUCCESS_OR_ELSE(ret, return);
 
 
132
 
 
 
133
        // Create a command queue
 
 
134
        cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
 
 
135
        CL_SUCCESS_OR_ELSE(ret, return);
 
 
136
        
 
 
137
        // Create program object
 
 
138
        cl_program program = clCreateProgramWithSource(context, 1, (const char**)&source_str, NULL, &ret);
 
 
139
        CL_SUCCESS_OR_ELSE(ret, return);
 
 
140
 
 
 
141
        // Compile & link the program
 
 
142
        ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
 
 
143
        //ret = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
 
 
144
        if (ret != CL_SUCCESS)
 
 
145
        {
 
 
146
            cout << "Could not build CL Program: " << GetCLErrorString(ret) << "\n";
 
 
147
            /*cl_build_status status;
 
 
148
            ret = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_STATUS, sizeof(status), &status, nullptr);
 
 
149
            assert(ret == CL_SUCCESS);*/
 
 
150
            //CL_SUCCESS_OR_ELSE(status, return);
 
 
151
            
 
 
152
            size_t length;
 
 
153
            static const int bufferSize = 8096;
 
 
154
            char buffer[bufferSize];
 
 
155
            ret = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, bufferSize, buffer, &length);
 
 
156
            assert(ret == CL_SUCCESS);
 
 
157
 
 
 
158
            buffer[length] = 0;
 
 
159
 
 
 
160
            cout << "--- Build log (" << length << " bytes) ---\n "<<buffer<<endl;
 
 
161
            return;
 
 
162
        }
 
 
163
 
 
 
164
        // Create the OpenCL kernel
 
 
165
        cl_kernel kernel = clCreateKernel(program, "solver", &ret);
 
 
166
        CL_SUCCESS_OR_ELSE(ret, return);
 
 
167
 
 
 
168
        // delete the source code
 
 
169
        delete [] source_str;
 
 
170
 
 
 
171
 
 
 
172
        // Prepare the data
 
 
173
        ParticleVector& particles = system->GetParticles();
 
 
174
        int nParticles = particles.size();
129
 
175
 
 
 
176
        const FluidVector& fluids = system->GetFluids();
 
 
177
        int nFluids = fluids.size();
 
 
178
        MaterialProperties materialProps;
 
 
179
        materialProps.masses = new Real[nFluids];
 
 
180
        materialProps.gasStiffness=  new Real[nFluids];
130
 
181
 
131
        Vector* positions = new Vector[nParticles];
182
        Vector* positions = new Vector[nParticles];
132
        Real* densities = new Real[nParticles];
183
        Real* densities = new Real[nParticles];
133
        int* materialIndexes = new int[nParticles];
184
        int* materialIndexes = new int[nParticles];
134
        
185
 
135
        //Particle neighbors[n_particles];
186
        //Particle neighbors[n_particles];
136
        int index = 0;
 
 
137
        for (int p = 0; p < nParticles; ++p)
187
        for (int p = 0; p < nParticles; ++p)
138
        {
188
        {
139
             positions[p] = particles[p].position;
189
            positions[p] = particles[p].position;
140
             densities[p] = particles[p].density;
190
            densities[p] = particles[p].density;
141
             materialIndexes[p] = MAT0;
191
            materialIndexes[p] = MAT0;
142
             index++;
192
        }
 
 
193
 
 
 
194
        for (int fluid = 0; fluid < nFluids; ++fluid)
 
 
195
        {
 
 
196
            materialProps.masses[fluid] = fluids[fluid]->GetMaterial()->GetMass();
 
 
197
            materialProps.gasStiffness[fluid] = fluids[fluid]->GetMaterial()->GetGasStiffness();
143
        }
198
        }
144
 
199
 
145
        
200
        // Allocate buffers on device
146
        int size_4_new_densities = nParticles * sizeof(Real);
201
        int sizePositions = sizeof(Vector) * nParticles;
147
        int size_4_pressure = nParticles * sizeof(Real);
202
        int sizeParticleScalar = sizeof(Real) * nParticles;
148
                
203
        int sizeMatIndices = sizeof(int) * nParticles;
149
        cl_mem d_positions = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(positions), NULL, &ret);
204
        int sizeMaterials = sizeof(MaterialProperties) * nFluids;
150
        cl_mem d_densities = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(densities), NULL, &ret);
205
 
151
        cl_mem d_materialIndexes = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(materialIndexes), NULL, &ret);
206
        cl_mem d_positions = clCreateBuffer(context, CL_MEM_READ_ONLY, sizePositions, NULL, &ret);
152
        cl_mem d_materials = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(materialParam), NULL, &ret);
207
        cl_mem d_densities = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeParticleScalar, NULL, &ret);
153
 
208
        cl_mem d_materialIndexes = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeMatIndices, NULL, &ret);
154
        cl_mem d_new_densities = clCreateBuffer(context, CL_MEM_READ_ONLY, size_4_new_densities, NULL, &ret);
209
        cl_mem d_materials = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeMaterials, NULL, &ret);
155
        cl_mem d_pressure = clCreateBuffer(context, CL_MEM_READ_ONLY, size_4_pressure, NULL, &ret);
210
 
156
 
211
        cl_mem d_new_densities = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeParticleScalar, NULL, &ret);
157
        // Create a command queue
212
        cl_mem d_pressure = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeParticleScalar, NULL, &ret);
158
        cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
213
 
159
 
214
        // Copy the data to their respective device buffers
160
        // Copy the lists to their respective memory buffers
215
        static const int blocking = CL_TRUE;
161
        ret = clEnqueueWriteBuffer(command_queue, d_positions, CL_TRUE, 0, sizeof(positions), positions, 0, NULL, NULL);
216
        ret = clEnqueueWriteBuffer(command_queue, d_positions, blocking, 0, sizePositions, positions, 0, NULL, NULL);
162
        ret = clEnqueueWriteBuffer(command_queue, d_densities, CL_TRUE, 0, sizeof(densities), densities, 0, NULL, NULL);
217
        ret = clEnqueueWriteBuffer(command_queue, d_densities, blocking, 0, sizeParticleScalar, densities, 0, NULL, NULL);
163
        ret = clEnqueueWriteBuffer(command_queue, d_materialIndexes, CL_TRUE, 0, sizeof(materialIndexes), materialIndexes, 0, NULL, NULL);
218
        ret = clEnqueueWriteBuffer(command_queue, d_materialIndexes, blocking, 0, sizeMatIndices, materialIndexes, 0, NULL, NULL);
164
        ret = clEnqueueWriteBuffer(command_queue, d_materials, CL_TRUE, 0, sizeof(materialParam), &materialParam, 0, NULL, NULL);
219
        ret = clEnqueueWriteBuffer(command_queue, d_materials, blocking, 0, sizeMaterials, &materialProps, 0, NULL, NULL);
165
        
220
 
166
        // Create a program from the kernel source
221
 
167
        cl_program program = clCreateProgramWithSource(context, 1, 
222
        // Set the arguments of the kernel
168
            (const char **)&source_str, (const size_t *)&source_size, &ret);
223
        ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_positions);
169
 
224
        ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&d_densities);
170
        //Build the program
225
        ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&d_materialIndexes);
171
        ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
226
        ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *)&d_materials);
172
 
227
 
173
        // Create the OpenCL kernel
228
        ret = clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *)&d_new_densities);
174
        cl_kernel kernel = clCreateKernel(program, "solver", &ret);
229
        ret = clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *)&d_pressure);
175
    
230
 
176
        // Set the arguments of the kernel
231
        ret = clSetKernelArg(kernel, 6, sizeof(nParticles), &nParticles);
177
        ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_positions);
232
 
178
        ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&d_densities);
233
 
179
        ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&d_materialIndexes);
234
        // Execute the OpenCL kernel on the list
180
        ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *)&d_materials);
235
        size_t local_work_size = 64; // Process one item at a time //Neither this
181
 
236
        size_t global_work_size = ((int)((int) local_work_size/ nParticles))+1;
182
        ret = clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *)&d_new_densities);
237
        ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
183
        ret = clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *)&d_pressure);
238
 
184
 
239
        // Read the memory buffer C on the device to the local variable C
185
        ret = clSetKernelArg(kernel, 6, sizeof(nParticles), &nParticles);
240
        Real* newDensities = new Real[nParticles];
186
 
241
        Real* newPressures = new Real[nParticles];
187
        
242
        ret = clEnqueueReadBuffer(command_queue, d_new_densities, CL_TRUE, 0, nParticles * sizeof(Real), newDensities, 0, NULL, NULL);
188
         // Execute the OpenCL kernel on the list
243
        ret = clEnqueueReadBuffer(command_queue, d_pressure, CL_TRUE, 0, nParticles * sizeof(Real), newPressures, 0, NULL, NULL);
189
        size_t local_work_size = 64; // Process one item at a time //Neither this
244
 
190
        size_t global_work_size = ((int)((int) local_work_size/ nParticles))+1;
 
 
191
        ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
 
 
192
 
 
 
193
        // Read the memory buffer C on the device to the local variable C
 
 
194
        Real* newDensities = new Real[nParticles];
 
 
195
        Real* newPressures = new Real[nParticles];
 
 
196
        ret = clEnqueueReadBuffer(command_queue, d_new_densities, CL_TRUE, 0, nParticles * sizeof(Real), newDensities, 0, NULL, NULL);
 
 
197
        ret = clEnqueueReadBuffer(command_queue, d_pressure, CL_TRUE, 0, nParticles * sizeof(Real), newPressures, 0, NULL, NULL);
 
 
198
        
 
 
199
        // Cleanup allocated objects
245
        // Cleanup allocated objects
200
        clReleaseKernel (kernel);  
246
        clReleaseKernel (kernel);  
201
        clReleaseProgram (program);
247
        clReleaseProgram (program);
202
        clReleaseCommandQueue (command_queue);
248
        clReleaseCommandQueue (command_queue);
203
        clReleaseContext (context);
249
        clReleaseContext (context);
204
        clReleaseMemObject (d_positions);
250
        clReleaseMemObject (d_positions);
205
        clReleaseMemObject (d_densities);
251
        clReleaseMemObject (d_densities);
206
        clReleaseMemObject (d_materialIndexes);
252
        clReleaseMemObject (d_materialIndexes);
207
        clReleaseMemObject (d_materials);
253
        clReleaseMemObject (d_materials);
208
        clReleaseMemObject (d_new_densities);
254
        clReleaseMemObject (d_new_densities);
209
        clReleaseMemObject (d_pressure);
255
        clReleaseMemObject (d_pressure);
210
        //free (cdDevices);
256
        //free (cdDevices);
211
        //free (cPathAndName);
257
        //free (cPathAndName);
212
        free (source_str);
258
        // Free host memory
213
        // Free host memory
259
        //free(srcA); 
214
        //free(srcA); 
260
        //free(srcB);
215
        //free(srcB);
261
        //free (dst);
216
        //free (dst);
262
 
217
        
263
        /*for (; _p != to; ++_p)
218
        /*for (; _p != to; ++_p)
 
 
219
        {
264
        {
220
            Particle& p = **_p;
265
        Particle& p = **_p;
221
 
266
 
222
            Real density = 0;
267
        Real density = 0;
223
            
 
 
224
            ParticlePtrVector neis = GetNeighbors(p);
 
 
225
 
268
 
226
            for (ParticlePtrIteratorConst neighbor = neis.begin(), end = neis.end(); neighbor != end; ++neighbor)
269
        ParticlePtrVector neis = GetNeighbors(p);
227
            {
270
        for (ParticlePtrIteratorConst neighbor = neis.begin(), end = neis.end(); neighbor != end; ++neighbor)
228
                Particle& nei = **neighbor;
271
        {
229
                density += gaussianKernel->evaluate(p.predictedPosition - nei.predictedPosition);
272
        Particle& nei = **neighbor;
230
            }
273
        density += gaussianKernel->evaluate(p.position - nei.position);
231
            density *= p.GetMass();
274
        }
 
 
275
        density *= p.GetMass();
 
 
276
        p.density = density;
232
 
277
 
233
            p.density = density;
278
        p.pressure = p.GetMaterial()->GetGasStiffness() * density;
234
            p.pressure = p.GetMaterial()->GetGasStiffness() * density;
 
 
235
        }*/
279
        }*/
236
    }
280
    }
237
 
281
 
...
 
...
 
265
            // Leapfrog integration
309
            // Leapfrog integration
266
            // see http://en.wikipedia.org/wiki/Leapfrog_integration
310
            // see http://en.wikipedia.org/wiki/Leapfrog_integration
267
 
311
 
268
            p.acceleration = p.GetSumOfAllForces() / p.GetMaterial()->GetRestDensity();
312
            p.acceleration = p.GetSumOfAllForces() / p.GetMaterial()->GetRestDensity();
269
 
313
 
270
            //newPos = p.position + p.velocity * GetConfig().GetDeltaT() + accel * GetConfig().GetDeltaT()HalfSquared;
314
            //newPos = p.position + p.velocity * GetConfig().GetDeltaT() + accel * GetConfig().GetDeltaT()HalfSquared;
271
            //newVelocity = p.velocity + GetConfig().GetDeltaT()Half * (accel + newAccel);
315
            //newVelocity = p.velocity + GetConfig().GetDeltaT()Half * (accel + newAccel);
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
43
        return 0;
43
        return 0;
44
#endif
44
#endif
45
    }
45
    }
 
 
46
 
 
 
47
    long ComputeFileSize(FILE* file)
 
 
48
    {
 
 
49
        long pos = ftell(file);
 
 
50
        // seek to end
 
 
51
        fseek(file, 0L, SEEK_END);
 
 
52
        // get size
 
 
53
        long size = ftell(file);
 
 
54
        // seek to last pos
 
 
55
        fseek(file, pos, SEEK_SET);
 
 
56
        return size;
 
 
57
    }
46
} }
58
} }
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
2
    WIN32 APPLICATION : Squishy.SPHDemo Project Overview
2
    WIN32 APPLICATION : Squishy.SPHDemo Project Overview
3
========================================================================
3
========================================================================
4
 
4
 
5
AppWizard has created this Squishy.SPHDemo application for you.
5
C++11 support:
 
 
6
clang: http://clang.llvm.org/cxx_status.html g++: http://gcc.gnu.org/projects/cxx0x.html micros~1: http://blogs.msdn.com/b/vcblog/archive/2011/09/12/10209291.aspx
6
 
7
 
7
This file contains a summary of what you will find in each of the files that
8
Also see: http://wiki.apache.org/stdcxx/C%2B%2B0xCompilerSupport
8
make up your Squishy.SPHDemo application.
 
 
9
 
 
 
10
 
 
 
11
Squishy.SPHDemo.vcxproj
 
 
12
    This is the main project file for VC++ projects generated using an Application Wizard.
 
 
13
    It contains information about the version of Visual C++ that generated the file, and
 
 
14
    information about the platforms, configurations, and project features selected with the
 
 
15
    Application Wizard.
 
 
16
 
 
 
17
Squishy.SPHDemo.vcxproj.filters
 
 
18
    This is the filters file for VC++ projects generated using an Application Wizard. 
 
 
19
    It contains information about the association between the files in your project 
 
 
20
    and the filters. This association is used in the IDE to show grouping of files with
 
 
21
    similar extensions under a specific node (for e.g. ".cpp" files are associated with the
 
 
22
    "Source Files" filter).
 
 
23
 
 
 
24
Squishy.SPHDemo.cpp
 
 
25
    This is the main application source file.
 
 
26
 
 
 
27
/////////////////////////////////////////////////////////////////////////////
 
 
28
AppWizard has created the following resources:
 
 
29
 
 
 
30
Squishy.SPHDemo.rc
 
 
31
    This is a listing of all of the Microsoft Windows resources that the
 
 
32
    program uses.  It includes the icons, bitmaps, and cursors that are stored
 
 
33
    in the RES subdirectory.  This file can be directly edited in Microsoft
 
 
34
    Visual C++.
 
 
35
 
 
 
36
Resource.h
 
 
37
    This is the standard header file, which defines new resource IDs.
 
 
38
    Microsoft Visual C++ reads and updates this file.
 
 
39
 
 
 
40
Squishy.SPHDemo.ico
 
 
41
    This is an icon file, which is used as the application's icon (32x32).
 
 
42
    This icon is included by the main resource file Squishy.SPHDemo.rc.
 
 
43
 
 
 
44
small.ico
 
 
45
    This is an icon file, which contains a smaller version (16x16)
 
 
46
    of the application's icon. This icon is included by the main resource
 
 
47
    file Squishy.SPHDemo.rc.
 
 
48
 
 
 
49
/////////////////////////////////////////////////////////////////////////////
 
 
50
Other standard files:
 
 
51
 
 
 
52
StdAfx.h, StdAfx.cpp
 
 
53
    These files are used to build a precompiled header (PCH) file
 
 
54
    named Squishy.SPHDemo.pch and a precompiled types file named StdAfx.obj.
 
 
55
 
 
 
56
/////////////////////////////////////////////////////////////////////////////
 
 
57
Other notes:
 
 
58
 
 
 
59
AppWizard uses "TODO:" comments to indicate parts of the source code you
 
 
60
should add to or customize.
 
 
61
 
 
 
62
/////////////////////////////////////////////////////////////////////////////
 
 
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
25
using namespace Squishy::SPH;
25
using namespace Squishy::SPH;
26
using namespace Squishy::SPH::Demo;
26
using namespace Squishy::SPH::Demo;
27
using namespace Squishy::SPH::Solvers;
27
using namespace Squishy::SPH::Solvers;
 
 
28
using namespace Squishy::SPH::DataStructures;
28
 
29
 
29
// TODO: Define world box
30
// TODO: Define world box
30
// TODO: Add simple render system for shapes
31
// TODO: Add simple render system for shapes
...
 
...
 
32
// ###################################################################################################
33
// ###################################################################################################
33
// Constants & Global variables
34
// Constants & Global variables
34
 
35
 
 
 
36
#define SolverCount 2
 
 
37
 
35
const int WindowWidth = 800;
38
const int WindowWidth = 800;
36
const int WindowHeight = 600;
39
const int WindowHeight = 600;
37
 
40
 
...
 
...
 
80
            SPHConfig cfg;
83
            SPHConfig cfg;
81
 
84
 
82
            shared_ptr<System> system = nullptr;
85
            shared_ptr<System> system = nullptr;
83
            shared_ptr<Solver> cpuSolver;
86
            shared_ptr<Solver> solvers[SolverCount];
84
            shared_ptr<Solver> gpuSolver;
87
            int solverIndex = 0;
85
 
88
 
86
            shared_ptr<Material> waterMaterial;
89
            shared_ptr<Material> waterMaterial;
87
            shared_ptr<Material> honeyMaterial;
90
            shared_ptr<Material> honeyMaterial;
...
 
...
 
103
                    // setup the solver(s) and the system(s)
106
                    // setup the solver(s) and the system(s)
104
                    cfg.HashtableResolution = 1;
107
                    cfg.HashtableResolution = 1;
105
 
108
 
106
                    cpuSolver = shared_ptr<Solver>(new WSPHSolver());
109
                    solvers[0] = shared_ptr<Solver>(new WSPHSolver());
107
                    gpuSolver = shared_ptr<Solver>(new WSPHGPUSolver());
110
                    solvers[1] = shared_ptr<Solver>(new WSPHGPUSolver());
108
 
111
 
109
                    system = shared_ptr<System>(new System(cfg, cpuSolver));
112
                    system = shared_ptr<System>(new System(cfg, solvers[solverIndex]));
110
 
113
 
111
                    waterMaterial = shared_ptr<Material>(new Water());
114
                    waterMaterial = shared_ptr<Material>(new Water());
112
                    honeyMaterial = shared_ptr<Material>(new Honey());
115
                    honeyMaterial = shared_ptr<Material>(new Honey());
...
 
...
 
201
                        }
204
                        }
202
                        else if (eventMask & EventMaskAddFluid)
205
                        else if (eventMask & EventMaskAddFluid)
203
                        {
206
                        {
 
 
207
                            // add some fluid
204
                            AddFluid(addParticleCount, addMaterial);
208
                            AddFluid(addParticleCount, addMaterial);
205
                        }
209
                        }
 
 
210
                        else if (eventMask & EventMaskSwitchSolver)
 
 
211
                        {
 
 
212
                            // switch to next solver
 
 
213
                            solverIndex = (solverIndex + 1) % SolverCount;
 
 
214
                            system->SetSolver(solvers[solverIndex]);
 
 
215
                        }
206
 
216
 
207
                        if (system->GetParticleCount() > 0)
217
                        if (system->GetParticleCount() > 0)
208
                        {
218
                        {
...
 
...
 
414
                case SDLK_r:
424
                case SDLK_r:
415
                    eventMask |= EventMaskReset;
425
                    eventMask |= EventMaskReset;
416
                    break;
426
                    break;
 
 
427
                case SDLK_o:
 
 
428
                    eventMask |= EventMaskSwitchSolver;
 
 
429
                    break;
417
 
430
 
418
                case SDLK_RETURN:
431
                case SDLK_RETURN:
419
                    eventMask |= EventMaskRun;
432
                    eventMask |= EventMaskRun;
...
 
...
 
793
                            // only draw relevant cells
806
                            // only draw relevant cells
794
                            if (
807
                            if (
795
                                (!sel && hashtable.IsOccupied(idx)) || 
808
                                (!sel && hashtable.IsOccupied(idx)) || 
796
                                (sel && hashtable.Hash(idx) == hashtable.Hash(sel->GetPosition())))
809
                                (sel && hashtable.MapIndex(idx) == hashtable.MapPos(sel->GetPosition())))
797
                            {
810
                            {
798
                                Vec3f current = Vec3f(idx) * l;
811
                                Vec3f current = Vec3f(idx) * l;
799
                                Vec3f minv = current + offsetMin;
812
                                Vec3f minv = current + offsetMin;
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
61
            EventMaskAddFluid =         0x0004,
61
            EventMaskAddFluid =         0x0004,
62
            EventMaskReset =            0x0008,
62
            EventMaskReset =            0x0008,
63
            EventMaskPickParticle =     0x0010,
63
            EventMaskPickParticle =     0x0010,
64
            EventMaskApplyForce =       0x0020
64
            EventMaskApplyForce =       0x0020,
 
 
65
            EventMaskSwitchSolver =     0x0040
65
    };
66
    };
66
 
67
 
67
    FLAGS(SystemEventMask)
68
    FLAGS(SystemEventMask)
b24387f1eb28ea2a3d4fc1681d68201948b4e82c7838344a88937312f99744af850f8b2fba4190c4
40
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
40
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
41
    <LinkIncremental>true</LinkIncremental>
41
    <LinkIncremental>true</LinkIncremental>
42
    <IncludePath>$(SolutionDir)include;$(ProjectDir)include;$(SolutionDir)include/glm;$(SolutionDir)Squishy.SPH/include;$(IncludePath)</IncludePath>
42
    <IncludePath>$(SolutionDir)include;$(ProjectDir)include;$(SolutionDir)include/glm;$(SolutionDir)Squishy.SPH/include;$(IncludePath)</IncludePath>
43
    <LibraryPath>$(SolutionDir)lib;$(LibraryPath)</LibraryPath>
43
    <LibraryPath>$(SolutionDir)lib;$(AMDAPPSDKROOT)\lib\x86;$(LibraryPath)</LibraryPath>
44
  </PropertyGroup>
44
  </PropertyGroup>
45
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
45
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
46
    <LinkIncremental>false</LinkIncremental>
46
    <LinkIncremental>false</LinkIncremental>
...
 
...
 
55
      <PreprocessorDefinitions>WIN32;_DEBUG;_WINDOWS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
55
      <PreprocessorDefinitions>WIN32;_DEBUG;_WINDOWS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
56
      <DisableSpecificWarnings>4819;4251;%(DisableSpecificWarnings)</DisableSpecificWarnings>
56
      <DisableSpecificWarnings>4819;4251;%(DisableSpecificWarnings)</DisableSpecificWarnings>
57
      <RuntimeLibrary>MultiThreadedDebugDLL</RuntimeLibrary>
57
      <RuntimeLibrary>MultiThreadedDebugDLL</RuntimeLibrary>
 
 
58
      <AdditionalIncludeDirectories>"$(AMDAPPSDKROOT)\include";</AdditionalIncludeDirectories>
58
    </ClCompile>
59
    </ClCompile>
59
    <Link>
60
    <Link>
60
      <SubSystem>Console</SubSystem>
61
      <SubSystem>Console</SubSystem>