• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@gmail.com>
3 //
4 // Distributed under the Boost Software License, Version 1.0
5 // See accompanying file LICENSE_1_0.txt or copy at
6 // http://www.boost.org/LICENSE_1_0.txt
7 //
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
10 
11 #include <iostream>
12 #include <algorithm>
13 
14 #include <QtGlobal>
15 #if QT_VERSION >= 0x050000
16 #include <QtWidgets>
17 #else
18 #include <QtGui>
19 #endif
20 #include <QtOpenGL>
21 
22 #ifndef Q_MOC_RUN
23 #include <boost/compute/command_queue.hpp>
24 #include <boost/compute/kernel.hpp>
25 #include <boost/compute/program.hpp>
26 #include <boost/compute/system.hpp>
27 #include <boost/compute/interop/opengl.hpp>
28 #include <boost/compute/utility/dim.hpp>
29 #include <boost/compute/utility/source.hpp>
30 #endif // Q_MOC_RUN
31 
32 namespace compute = boost::compute;
33 
34 // opencl source code
35 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
36     // map value to color
37     float4 color(uint i)
38     {
39         uchar c = i;
40         uchar x = 35;
41         uchar y = 25;
42         uchar z = 15;
43         uchar max = 255;
44 
45         if(i == 256)
46             return (float4)(0, 0, 0, 255);
47         else
48             return (float4)(max-x*i, max-y*i, max-z*i, max) / 255.f;
49     }
50 
51     __kernel void mandelbrot(__write_only image2d_t image)
52     {
53         const uint x_coord = get_global_id(0);
54         const uint y_coord = get_global_id(1);
55         const uint width = get_global_size(0);
56         const uint height = get_global_size(1);
57 
58         float x_origin = ((float) x_coord / width) * 3.25f - 2.0f;
59         float y_origin = ((float) y_coord / height) * 2.5f - 1.25f;
60 
61         float x = 0.0f;
62         float y = 0.0f;
63 
64         uint i = 0;
65         while(x*x + y*y <= 4.f && i < 256){
66             float tmp = x*x - y*y + x_origin;
67             y = 2*x*y + y_origin;
68             x = tmp;
69             i++;
70         }
71 
72         int2 coord = { x_coord, y_coord };
73         write_imagef(image, coord, color(i));
74     };
75 );
76 
77 class MandelbrotWidget : public QGLWidget
78 {
79     Q_OBJECT
80 
81 public:
82     MandelbrotWidget(QWidget *parent = 0);
83     ~MandelbrotWidget();
84 
85     void initializeGL();
86     void resizeGL(int width, int height);
87     void paintGL();
88     void keyPressEvent(QKeyEvent* event);
89 
90 private:
91     compute::context context_;
92     compute::command_queue queue_;
93     compute::program program_;
94     GLuint gl_texture_;
95     compute::opengl_texture cl_texture_;
96 };
97 
MandelbrotWidget(QWidget * parent)98 MandelbrotWidget::MandelbrotWidget(QWidget *parent)
99     : QGLWidget(parent)
100 {
101     gl_texture_ = 0;
102 }
103 
~MandelbrotWidget()104 MandelbrotWidget::~MandelbrotWidget()
105 {
106 }
107 
initializeGL()108 void MandelbrotWidget::initializeGL()
109 {
110     // setup opengl
111     glDisable(GL_LIGHTING);
112 
113     // create the OpenGL/OpenCL shared context
114     context_ = compute::opengl_create_shared_context();
115 
116     // get gpu device
117     compute::device gpu = context_.get_device();
118     std::cout << "device: " << gpu.name() << std::endl;
119 
120     // setup command queue
121     queue_ = compute::command_queue(context_, gpu);
122 
123     // build mandelbrot program
124     program_ = compute::program::create_with_source(source, context_);
125     program_.build();
126 }
127 
resizeGL(int width,int height)128 void MandelbrotWidget::resizeGL(int width, int height)
129 {
130 #if QT_VERSION >= 0x050000
131     // scale height/width based on device pixel ratio
132     width /= windowHandle()->devicePixelRatio();
133     height /= windowHandle()->devicePixelRatio();
134 #endif
135 
136     // resize viewport
137     glViewport(0, 0, width, height);
138 
139     // delete old texture
140     if(gl_texture_){
141       glDeleteTextures(1, &gl_texture_);
142       gl_texture_ = 0;
143     }
144 
145     // generate new texture
146     glGenTextures(1, &gl_texture_);
147     glBindTexture(GL_TEXTURE_2D, gl_texture_);
148     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_REPEAT);
149     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_REPEAT);
150     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
151     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
152     glTexImage2D(
153         GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0
154     );
155 
156     // create opencl object for the texture
157     cl_texture_ = compute::opengl_texture(
158         context_, GL_TEXTURE_2D, 0, gl_texture_, CL_MEM_WRITE_ONLY
159     );
160 }
161 
paintGL()162 void MandelbrotWidget::paintGL()
163 {
164     using compute::dim;
165 
166     float w = width();
167     float h = height();
168 
169     glMatrixMode(GL_PROJECTION);
170     glLoadIdentity();
171     glOrtho(0.0, w, 0.0, h, -1.0, 1.0);
172     glMatrixMode(GL_MODELVIEW);
173     glLoadIdentity();
174 
175     // setup the mandelbrot kernel
176     compute::kernel kernel(program_, "mandelbrot");
177     kernel.set_arg(0, cl_texture_);
178 
179     // acquire the opengl texture so it can be used in opencl
180     compute::opengl_enqueue_acquire_gl_objects(1, &cl_texture_.get(), queue_);
181 
182     // execute the mandelbrot kernel
183     queue_.enqueue_nd_range_kernel(
184         kernel, dim(0, 0), dim(width(), height()), dim(1, 1)
185     );
186 
187     // release the opengl texture so it can be used by opengl
188     compute::opengl_enqueue_release_gl_objects(1, &cl_texture_.get(), queue_);
189 
190     // ensure opencl is finished before rendering in opengl
191     queue_.finish();
192 
193     // draw a single quad with the mandelbrot image texture
194     glEnable(GL_TEXTURE_2D);
195     glBindTexture(GL_TEXTURE_2D, gl_texture_);
196 
197     glBegin(GL_QUADS);
198     glTexCoord2f(0, 0); glVertex2f(0, 0);
199     glTexCoord2f(0, 1); glVertex2f(0, h);
200     glTexCoord2f(1, 1); glVertex2f(w, h);
201     glTexCoord2f(1, 0); glVertex2f(w, 0);
202     glEnd();
203 }
204 
keyPressEvent(QKeyEvent * event)205 void MandelbrotWidget::keyPressEvent(QKeyEvent* event)
206 {
207     if(event->key() == Qt::Key_Escape) {
208         this->close();
209     }
210 }
211 
212 // the mandelbrot example shows how to create a mandelbrot image in
213 // OpenCL and render the image as a texture in OpenGL
main(int argc,char * argv[])214 int main(int argc, char *argv[])
215 {
216     QApplication app(argc, argv);
217 
218     MandelbrotWidget widget;
219     widget.show();
220 
221     return app.exec();
222 }
223 
224 #include "mandelbrot.moc"
225