From 878e7d5cf27a36e4bba8b8140311fb0135bb29cb Mon Sep 17 00:00:00 2001 From: Joshua Moerman Date: Tue, 29 Apr 2014 11:50:47 +0200 Subject: [PATCH] Uses opencl to generate a texture --- resources/Fractal.fsh | 21 +----- resources/Fractal.vsh | 15 ++--- resources/Kernel.cl | 18 ++++-- src/main.cpp | 144 ++++++++++++++++++++++++++++++------------ 4 files changed, 123 insertions(+), 75 deletions(-) diff --git a/resources/Fractal.fsh b/resources/Fractal.fsh index d4ca0da..d0b5677 100644 --- a/resources/Fractal.fsh +++ b/resources/Fractal.fsh @@ -1,26 +1,11 @@ #version 330 -uniform float rotation; +uniform sampler2D tex; -in vec2 pos; -in vec2 start; +in vec2 tex_coord_int; out vec4 fragColor; void main(){ - - vec2 z = start; - - int i = 0; - for (i = 0; i < 30; i++) { - vec2 zsq = z*z; - if(zsq.x + zsq.y > 16.0) break; - - float t = zsq.x - zsq.y + pos.x; - z.y = 2.0*z.x*z.y + pos.y; - z.x = t; - } - - fragColor = vec4(float(i) / 30.0); - fragColor.bg = 0.5 * sin(z) + 0.5; + fragColor = texture(tex, tex_coord_int); } diff --git a/resources/Fractal.vsh b/resources/Fractal.vsh index 329f2d2..cc0716c 100644 --- a/resources/Fractal.vsh +++ b/resources/Fractal.vsh @@ -1,18 +1,11 @@ #version 330 -uniform float rotation; - in vec4 position; -in vec4 color; +in vec2 tex_coord; -out vec2 pos; -out vec2 start; +out vec2 tex_coord_int; void main(){ - pos = position.xy; - pos.x -= 0.5; - - start = sqrt(1.0 + 0.01 * rotation) * 0.3 * sin(vec2(0.1, 0.1337) * rotation); - - gl_Position = position; + tex_coord_int = tex_coord; + gl_Position = position; } diff --git a/resources/Kernel.cl b/resources/Kernel.cl index 8fc6628..e22cd69 100644 --- a/resources/Kernel.cl +++ b/resources/Kernel.cl @@ -1,8 +1,18 @@ +kernel void initialize(size_t width, size_t height, write_only image2d_t output){ + size_t x = get_global_id(0); + size_t y = get_global_id(1); + + int2 coord = {x, y}; + uint4 color = {255 * x / (width - 1), 255 * y / (height - 1), 0, 255}; + + write_imageui(output, coord, color); +} + kernel void square(global float* input, size_t width, global float* output){ - size_t x = get_global_id(0); - size_t y = get_global_id(1); - float i = 2.0 * (input[x + width*y] - 0.5); - output[x + width*y] = 0.5 + 0.5*sin((1.0 - i) * sin(i*i + 6.2*float(x)/width) + cos(1.0/i + 6.2*float(y)/width)); + size_t x = get_global_id(0); + size_t y = get_global_id(1); + float i = 2.0 * (input[x + width*y] - 0.5); + output[x + width*y] = 0.5 + 0.5*sin((1.0 - i) * sin(i*i + 6.2*float(x)/width) + cos(1.0/i + 6.2*float(y)/width)); } diff --git a/src/main.cpp b/src/main.cpp index 6cc90ae..7aad22b 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,6 +1,6 @@ #include - #include +#include #include #include @@ -8,62 +8,57 @@ #include #include -#include +#include +#include +#include #include #include #include -#include #include -#include +#include using namespace std; -#define BUFFER_OFFSET(i) ((char *)NULL + (i)) +std::string slurp(std::string file_name){ + std::ifstream f(file_name); + if (!f) throw std::runtime_error(std::string("Unable to open file: ") + file_name); + return {std::istreambuf_iterator(f), std::istreambuf_iterator()}; +} struct Vertex{ moggle::vector3 position; - moggle::vector3 color; + moggle::vector2 tex_coord; }; const Vertex quad[] = { - // x y z, r g b - {{1, -1, 0}, {1, 0, 0}}, - {{-1, -1, 0}, {0, 1, 0}}, - {{1, 1, 0}, {0, 0, 1}}, - {{-1, 1, 0}, {1, 1, 1}} + // x y z u v + {{ 1, -1, 0}, {1, 0}}, + {{-1, -1, 0}, {0, 0}}, + {{ 1, 1, 0}, {1, 1}}, + {{-1, 1, 0}, {0, 1}} }; -static void check_shader(GLuint s){ - GLint logLength; - glGetShaderiv(s, GL_INFO_LOG_LENGTH, &logLength); - if (logLength > 0) { - GLchar *log = (GLchar *)malloc(logLength); - glGetShaderInfoLog(s, logLength, &logLength, log); - std::cout << "Shader compile log:\n%s" << log << std::endl; - free(log); - } -} - -string slurp(string filename) { - ifstream in(filename); - stringstream sstr; - sstr << in.rdbuf(); - return sstr.str(); -} - struct App { - float time = 0; moggle::shader_program program; moggle::vao v; moggle::vbo quad_vbo; + // ptrs to postpone construction, should clean this up + unique_ptr cl_context; + unique_ptr cl_image; + + float time = 0; + GLuint gl_texture = 0; + size_t W = 128; + size_t H = 128; + enum { POS_ATTR, - COL_ATTR + TEX_ATTR }; - void initialize(){ + void load_shader(){ auto vs = moggle::shader::from_file(moggle::shader_type::vertex, "Fractal.vsh"); auto fs = moggle::shader::from_file(moggle::shader_type::fragment, "Fractal.fsh"); @@ -71,19 +66,81 @@ struct App { program.attach(fs); program.bind_attribute(POS_ATTR, "position"); - program.bind_attribute(COL_ATTR, "color"); + program.bind_attribute(TEX_ATTR, "tex_coord"); program.link(); + } + void create_vao(){ v.bind(); quad_vbo.bind(GL_ARRAY_BUFFER); quad_vbo.data(quad); moggle::gl::enable_vertex_attribute_array(POS_ATTR); - moggle::gl::enable_vertex_attribute_array(COL_ATTR); + moggle::gl::enable_vertex_attribute_array(TEX_ATTR); v.attribute(POS_ATTR, quad_vbo, &Vertex::position); - v.attribute(COL_ATTR, quad_vbo, &Vertex::color); + v.attribute(TEX_ATTR, quad_vbo, &Vertex::tex_coord); + } + + GLuint create_texture(){ + std::vector data(4*W*H); + + for(int i = 0; i < data.size(); i += 4){ + data[i+0] = 255 * i / (W-1); + data[i+1] = 127; + data[i+2] = 0; + data[i+3] = 255; + } + + GLuint tex = 0; + moggle::gl::active_texture(GL_TEXTURE0); + moggle::gl::generate_textures(1, &tex); + moggle::gl::bind_texture(GL_TEXTURE_2D, tex); + moggle::gl::texture_image_2d(GL_TEXTURE_2D, 0, GL_RGBA, W, H, 0, GL_RGBA, GL_UNSIGNED_BYTE, data.data()); + moggle::gl::texture_parameter_i(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + moggle::gl::texture_parameter_i(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + return tex; + } + + void initialize(GLContext & gl_context){ + // Fractal part + load_shader(); + create_vao(); + + // create a texture for opencl + gl_texture = create_texture(); + + // share the bitch + CGLShareGroupObj sharegroup = CGLGetShareGroup(gl_context.ctx); + + cl_int err = 0; + cl_context_properties properties[] = { CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, (cl_context_properties)sharegroup, 0 }; + + cl_context = make_unique(CL_DEVICE_TYPE_CPU, properties, nullptr, nullptr, &err); + cl::checky(err); + auto cl_device = cl_context->getInfo().front(); + + // obtain cl memory object + cl_image = make_unique(*cl_context, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, gl_texture, &err); + cl::checky(err); + std::cout << *cl_image << std::endl; // sanity check + + // build the program + auto KernelSource = slurp("Kernel.cl"); + cl::Program program(*cl_context, {KernelSource.c_str(), KernelSource.size()}, true, &err); + cout << program.getBuildInfo(cl_device) << std::endl; + cl::checky(err); + cout << program << endl; + + // grab the kernel + KernelOp kernel(program, "initialize", &err); + cl::checky(err); + + // create a queue (for the first device) + cl::CommandQueue queue(*cl_context, cl_device); + cl::checky(kernel(queue, W, H, W, H, *cl_image)); + cl::checky(queue.flush()); } void draw(){ @@ -93,7 +150,10 @@ struct App { moggle::gl::clear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); program.use(); - program.uniform("rotation").set(time); + + moggle::gl::bind_texture(GL_TEXTURE_2D, gl_texture); + program.uniform("tex").set(0); + v.bind(); moggle::gl::draw_arrays(GL_TRIANGLE_STRIP, 0, 4); @@ -120,8 +180,8 @@ struct CLApp { cl::CommandQueue queue(context, context.getInfo().front()); // make a lot of data - constexpr size_t W = 1280 * 1; - constexpr size_t H = 800 * 1; + constexpr size_t W = 512; + constexpr size_t H = 512; std::vector input_vector(W*H); for(int y = 0; y < H; ++y){ @@ -133,7 +193,7 @@ struct CLApp { // transfer data into buffers cl::Buffer input(context, input_vector.begin(), input_vector.end(), false, true); - int r = 20, g = 20, b = 20; + int r = 10, g = 10, b = 10; // DO IT (in place) for(int i = 0; i < r; ++i){ @@ -183,9 +243,9 @@ int main() { NSAppWrapper app; app.create_window({ - [&](ContextParameters){ - a.initialize(); - // b.initialize(); + [&](ContextParameters ctxp){ + //b.initialize(); + a.initialize(ctxp.context); }, [&](ContextParameters){ a.draw();