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"
|