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
|
//---------------------------------------------------------------------------//
// Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@gmail.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>
#include <algorithm>
#include <QtGlobal>
#if QT_VERSION >= 0x050000
#include <QtWidgets>
#else
#include <QtGui>
#endif
#include <QtOpenGL>
#ifndef Q_MOC_RUN
#include <boost/compute/command_queue.hpp>
#include <boost/compute/kernel.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/interop/opengl.hpp>
#include <boost/compute/utility/dim.hpp>
#include <boost/compute/utility/source.hpp>
#endif // Q_MOC_RUN
namespace compute = boost::compute;
// opencl source code
const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
// map value to color
float4 color(uint i)
{
uchar c = i;
uchar x = 35;
uchar y = 25;
uchar z = 15;
uchar max = 255;
if(i == 256)
return (float4)(0, 0, 0, 255);
else
return (float4)(max-x*i, max-y*i, max-z*i, max) / 255.f;
}
__kernel void mandelbrot(__write_only image2d_t image)
{
const uint x_coord = get_global_id(0);
const uint y_coord = get_global_id(1);
const uint width = get_global_size(0);
const uint height = get_global_size(1);
float x_origin = ((float) x_coord / width) * 3.25f - 2.0f;
float y_origin = ((float) y_coord / height) * 2.5f - 1.25f;
float x = 0.0f;
float y = 0.0f;
uint i = 0;
while(x*x + y*y <= 4.f && i < 256){
float tmp = x*x - y*y + x_origin;
y = 2*x*y + y_origin;
x = tmp;
i++;
}
int2 coord = { x_coord, y_coord };
write_imagef(image, coord, color(i));
};
);
class MandelbrotWidget : public QGLWidget
{
Q_OBJECT
public:
MandelbrotWidget(QWidget *parent = 0);
~MandelbrotWidget();
void initializeGL();
void resizeGL(int width, int height);
void paintGL();
void keyPressEvent(QKeyEvent* event);
private:
compute::context context_;
compute::command_queue queue_;
compute::program program_;
GLuint gl_texture_;
compute::opengl_texture cl_texture_;
};
MandelbrotWidget::MandelbrotWidget(QWidget *parent)
: QGLWidget(parent)
{
gl_texture_ = 0;
}
MandelbrotWidget::~MandelbrotWidget()
{
}
void MandelbrotWidget::initializeGL()
{
// setup opengl
glDisable(GL_LIGHTING);
// create the OpenGL/OpenCL shared context
context_ = compute::opengl_create_shared_context();
// get gpu device
compute::device gpu = context_.get_device();
std::cout << "device: " << gpu.name() << std::endl;
// setup command queue
queue_ = compute::command_queue(context_, gpu);
// build mandelbrot program
program_ = compute::program::create_with_source(source, context_);
program_.build();
}
void MandelbrotWidget::resizeGL(int width, int height)
{
#if QT_VERSION >= 0x050000
// scale height/width based on device pixel ratio
width /= windowHandle()->devicePixelRatio();
height /= windowHandle()->devicePixelRatio();
#endif
// resize viewport
glViewport(0, 0, width, height);
// delete old texture
if(gl_texture_){
glDeleteTextures(1, &gl_texture_);
gl_texture_ = 0;
}
// generate new texture
glGenTextures(1, &gl_texture_);
glBindTexture(GL_TEXTURE_2D, gl_texture_);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_REPEAT);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_REPEAT);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
glTexImage2D(
GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0
);
// create opencl object for the texture
cl_texture_ = compute::opengl_texture(
context_, GL_TEXTURE_2D, 0, gl_texture_, CL_MEM_WRITE_ONLY
);
}
void MandelbrotWidget::paintGL()
{
using compute::dim;
float w = width();
float h = height();
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glOrtho(0.0, w, 0.0, h, -1.0, 1.0);
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
// setup the mandelbrot kernel
compute::kernel kernel(program_, "mandelbrot");
kernel.set_arg(0, cl_texture_);
// acquire the opengl texture so it can be used in opencl
compute::opengl_enqueue_acquire_gl_objects(1, &cl_texture_.get(), queue_);
// execute the mandelbrot kernel
queue_.enqueue_nd_range_kernel(
kernel, dim(0, 0), dim(width(), height()), dim(1, 1)
);
// release the opengl texture so it can be used by opengl
compute::opengl_enqueue_release_gl_objects(1, &cl_texture_.get(), queue_);
// ensure opencl is finished before rendering in opengl
queue_.finish();
// draw a single quad with the mandelbrot image texture
glEnable(GL_TEXTURE_2D);
glBindTexture(GL_TEXTURE_2D, gl_texture_);
glBegin(GL_QUADS);
glTexCoord2f(0, 0); glVertex2f(0, 0);
glTexCoord2f(0, 1); glVertex2f(0, h);
glTexCoord2f(1, 1); glVertex2f(w, h);
glTexCoord2f(1, 0); glVertex2f(w, 0);
glEnd();
}
void MandelbrotWidget::keyPressEvent(QKeyEvent* event)
{
if(event->key() == Qt::Key_Escape) {
this->close();
}
}
// the mandelbrot example shows how to create a mandelbrot image in
// OpenCL and render the image as a texture in OpenGL
int main(int argc, char *argv[])
{
QApplication app(argc, argv);
MandelbrotWidget widget;
widget.show();
return app.exec();
}
#include "mandelbrot.moc"
|
