1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
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
77
78
79
80
81
82
83
84
85
86
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
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
|
//---------------------------------------------------------------------------//
// Copyright (c) 2014 Fabian Köhler <fabian2804@googlemail.com>
//
// Distributed under the Boost Software License, Version 1.0
// See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt
//
// See http://boostorg.github.com/compute for more information.
//---------------------------------------------------------------------------//
#include <iostream>
#define GL_GLEXT_PROTOTYPES
#ifdef __APPLE__
#include <OpenGL/gl.h>
#include <OpenGL/glext.h>
#else
#include <GL/gl.h>
#include <GL/glext.h>
#endif
#include <QtGlobal>
#if QT_VERSION >= 0x050000
#include <QtWidgets>
#else
#include <QtGui>
#endif
#include <QtOpenGL>
#include <QTimer>
#include <boost/program_options.hpp>
#include <boost/random/uniform_real_distribution.hpp>
#include <boost/random/mersenne_twister.hpp>
#ifndef Q_MOC_RUN
#include <boost/compute/system.hpp>
#include <boost/compute/algorithm.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/interop/opengl.hpp>
#include <boost/compute/utility/source.hpp>
#endif // Q_MOC_RUN
namespace compute = boost::compute;
namespace po = boost::program_options;
using compute::uint_;
using compute::float4_;
const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
__kernel void updateVelocity(__global const float4* position, __global float4* velocity, float dt, uint N)
{
uint gid = get_global_id(0);
float4 r = { 0.0f, 0.0f, 0.0f, 0.0f };
float f = 0.0f;
for(uint i = 0; i != gid; i++) {
if(i != gid) {
r = position[i]-position[gid];
f = length(r)+0.001f;
f *= f*f;
f = dt/f;
velocity[gid] += f*r;
}
}
}
__kernel void updatePosition(__global float4* position, __global const float4* velocity, float dt)
{
uint gid = get_global_id(0);
position[gid].xyz += dt*velocity[gid].xyz;
}
);
class NBodyWidget : public QGLWidget
{
Q_OBJECT
public:
NBodyWidget(std::size_t particles, float dt, QWidget* parent = 0);
~NBodyWidget();
void initializeGL();
void resizeGL(int width, int height);
void paintGL();
void updateParticles();
void keyPressEvent(QKeyEvent* event);
private:
QTimer* timer;
compute::context m_context;
compute::command_queue m_queue;
compute::program m_program;
compute::opengl_buffer m_position;
compute::vector<float4_>* m_velocity;
compute::kernel m_velocity_kernel;
compute::kernel m_position_kernel;
bool m_initial_draw;
const uint_ m_particles;
const float m_dt;
};
NBodyWidget::NBodyWidget(std::size_t particles, float dt, QWidget* parent)
: QGLWidget(parent), m_initial_draw(true), m_particles(particles), m_dt(dt)
{
// create a timer to redraw as fast as possible
timer = new QTimer(this);
connect(timer, SIGNAL(timeout()), this, SLOT(updateGL()));
timer->start(1);
}
NBodyWidget::~NBodyWidget()
{
delete m_velocity;
// delete the opengl buffer
GLuint vbo = m_position.get_opengl_object();
glDeleteBuffers(1, &vbo);
}
void NBodyWidget::initializeGL()
{
// create context, command queue and program
m_context = compute::opengl_create_shared_context();
m_queue = compute::command_queue(m_context, m_context.get_device());
m_program = compute::program::create_with_source(source, m_context);
m_program.build();
// prepare random particle positions that will be transferred to the vbo
float4_* temp = new float4_[m_particles];
boost::random::uniform_real_distribution<float> dist(-0.5f, 0.5f);
boost::random::mt19937_64 gen;
for(size_t i = 0; i < m_particles; i++) {
temp[i][0] = dist(gen);
temp[i][1] = dist(gen);
temp[i][2] = dist(gen);
temp[i][3] = 1.0f;
}
// create an OpenGL vbo
GLuint vbo = 0;
glGenBuffers(1, &vbo);
glBindBuffer(GL_ARRAY_BUFFER, vbo);
glBufferData(GL_ARRAY_BUFFER, m_particles*sizeof(float4_), temp, GL_DYNAMIC_DRAW);
// create a OpenCL buffer from the vbo
m_position = compute::opengl_buffer(m_context, vbo);
delete[] temp;
// create buffer for velocities
m_velocity = new compute::vector<float4_>(m_particles, m_context);
compute::fill(m_velocity->begin(), m_velocity->end(), float4_(0.0f, 0.0f, 0.0f, 0.0f), m_queue);
// create compute kernels
m_velocity_kernel = m_program.create_kernel("updateVelocity");
m_velocity_kernel.set_arg(0, m_position);
m_velocity_kernel.set_arg(1, m_velocity->get_buffer());
m_velocity_kernel.set_arg(2, m_dt);
m_velocity_kernel.set_arg(3, m_particles);
m_position_kernel = m_program.create_kernel("updatePosition");
m_position_kernel.set_arg(0, m_position);
m_position_kernel.set_arg(1, m_velocity->get_buffer());
m_position_kernel.set_arg(2, m_dt);
}
void NBodyWidget::resizeGL(int width, int height)
{
// update viewport
glViewport(0, 0, width, height);
}
void NBodyWidget::paintGL()
{
// clear buffer
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
// check if this is the first draw
if(m_initial_draw) {
// do not update particles
m_initial_draw = false;
} else {
// update particles
updateParticles();
}
// draw
glVertexPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_VERTEX_ARRAY);
glDrawArrays(GL_POINTS, 0, m_particles);
glFinish();
}
void NBodyWidget::updateParticles()
{
// enqueue kernels to update particles and make sure that the command queue is finished
compute::opengl_enqueue_acquire_buffer(m_position, m_queue);
m_queue.enqueue_1d_range_kernel(m_velocity_kernel, 0, m_particles, 0).wait();
m_queue.enqueue_1d_range_kernel(m_position_kernel, 0, m_particles, 0).wait();
m_queue.finish();
compute::opengl_enqueue_release_buffer(m_position, m_queue);
}
void NBodyWidget::keyPressEvent(QKeyEvent* event)
{
if(event->key() == Qt::Key_Escape) {
this->close();
}
}
int main(int argc, char** argv)
{
// parse command line arguments
po::options_description options("options");
options.add_options()
("help", "show usage")
("particles", po::value<uint_>()->default_value(1000), "number of particles")
("dt", po::value<float>()->default_value(0.00001f), "width of each integration step");
po::variables_map vm;
po::store(po::parse_command_line(argc, argv, options), vm);
po::notify(vm);
if(vm.count("help") > 0) {
std::cout << options << std::endl;
return 0;
}
const uint_ particles = vm["particles"].as<uint_>();
const float dt = vm["dt"].as<float>();
QApplication app(argc, argv);
NBodyWidget nbody(particles, dt);
nbody.show();
return app.exec();
}
#include "nbody.moc"
|