diff --git a/CMakeLists.txt b/CMakeLists.txt index abcb076..6962296 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -4,6 +4,7 @@ cmake_minimum_required(VERSION 2.8) add_definitions(-std=c++1y) add_subdirectory("contrib") +add_subdirectory("lib") add_subdirectory("src") file(GLOB resources "resources/*") diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt new file mode 100644 index 0000000..794f022 --- /dev/null +++ b/lib/CMakeLists.txt @@ -0,0 +1,9 @@ + +file(GLOB sources "*.cpp") +file(GLOB headers "*.hpp") + +set(libs common ${JCLWrapper_LIBRARIES} ${JNSAppWrapper_LIBRARIES} ${ImageStreams_LIBRARIES} moggle_xxx) + +add_library(common ${headers} ${sources}) +target_link_libraries(common ${libs}) +target_include_directories(common PUBLIC ".") diff --git a/lib/app.cpp b/lib/app.cpp new file mode 100644 index 0000000..9f46097 --- /dev/null +++ b/lib/app.cpp @@ -0,0 +1,174 @@ +#include "app.hpp" + +#include "utils.hpp" + +#include + +using namespace std; + +static const std::vector quad = {{ + // 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}} +}}; + +enum { + POS_ATTR, + TEX_ATTR +}; + +App::App(GLContext& gl_context) +: gl_image1(create_gl_texture(W, H)) +, gl_image2(create_gl_texture(W, H)) +, gl_vbo(create_vbo(quad)) +, gl_vao(create_vao(gl_vbo)) +, gl_program(create_shader_from_files("Fractal.vsh", "Fractal.fsh")) +, cl_context(create_shared_cl_context(gl_context)) +, cl_device(cl_context.getInfo().front()) +, cl_queue(cl_context, cl_device) +, cl_image1(cl_context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, gl_image1) +, cl_image2(cl_context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, gl_image2) +, cl_program(create_cl_program("Kernel.cl")) +, k_initialize(cl_program, "initialize") +, k_hblur(cl_program, "hblur") +, k_vblur(cl_program, "vblur") +, k_update(cl_program, "update") +{ + cl::checky(k_initialize(cl_queue, W, H, W, H, cl_image1)); + cl::checky(k_initialize(cl_queue, W, H, W, H, cl_image2)); +} + +void App::draw(){ + GLuint* gl_images[] = {&gl_image1, &gl_image2}; + cl::ImageGL* cl_images[] = {&cl_image1, &cl_image2}; + + auto dt = 1/60.0; + time += dt; + if(wait > 0) wait -= dt; + + if(wait <= 0){ + bool blur = false; + if((frame / 200) % 2) blur = true; + + if(blur){ + frame += 4; + cl::checky(k_hblur(cl_queue, W, H, W, H, *cl_images[(frame + 1) % 2], *cl_images[frame % 2])); + cl::checky(k_vblur(cl_queue, W, H, W, H, *cl_images[frame % 2], *cl_images[(frame + 1) % 2])); + } else { + frame++; + cl::checky(k_update(cl_queue, W, H, W, H, *cl_images[(frame + 1) % 2], *cl_images[frame % 2])); + } + } + + cl::checky(cl_queue.finish()); + + moggle::gl::clear_color(0.35f, 0.65f, 0.65f, 1.0f); + moggle::gl::clear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + + gl_program.use(); + + moggle::gl::bind_texture(GL_TEXTURE_2D, *gl_images[frame % 2]); + gl_program.uniform("tex").set(0); + + gl_vao.bind(); + + moggle::gl::draw_arrays(GL_TRIANGLE_STRIP, 0, 4); +} + +void App::resize(size_t w, size_t h){ + wait = 2; + + W = w; + H = h; + + gl_image1 = create_gl_texture(W, H); + gl_image2 = create_gl_texture(W, H); + + cl_int err = 0; + cl_image1 = cl::ImageGL(cl_context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, gl_image1, &err); + cl::checky(err); + cl_image2 = cl::ImageGL(cl_context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, gl_image2, &err); + cl::checky(err); + std::cout << cl_image1 << std::endl; + + cl::checky(k_initialize(cl_queue, W, H, W, H, cl_image1)); + cl::checky(k_initialize(cl_queue, W, H, W, H, cl_image2)); +} + +moggle::shader_program App::create_shader_from_files(string vertex, string fragment){ + moggle::shader_program program; + + auto vs = moggle::shader::from_file(moggle::shader_type::vertex, vertex); + auto fs = moggle::shader::from_file(moggle::shader_type::fragment, fragment); + + program.attach(vs); + program.attach(fs); + + program.bind_attribute(POS_ATTR, "position"); + program.bind_attribute(TEX_ATTR, "tex_coord"); + + program.link(); + + return program; +} + +moggle::vbo App::create_vbo(const std::vector& data){ + moggle::vbo vbo; + vbo.bind(GL_ARRAY_BUFFER); + vbo.data(data); + return vbo; +} + +moggle::vao App::create_vao(const moggle::vbo& vbo){ + moggle::vao vao; + vao.bind(); + + moggle::gl::enable_vertex_attribute_array(POS_ATTR); + moggle::gl::enable_vertex_attribute_array(TEX_ATTR); + + vao.attribute(POS_ATTR, vbo, &Vertex::position); + vao.attribute(TEX_ATTR, vbo, &Vertex::tex_coord); + + return vao; +} + +GLuint App::create_gl_texture(GLsizei width, GLsizei height){ + 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_RGBA32F, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, nullptr); + 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; +} + +cl::Context App::create_shared_cl_context(GLContext& gl_context){ + CGLShareGroupObj sharegroup = CGLGetShareGroup(gl_context.ctx); + + cl_context_properties properties[] = { + CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, (cl_context_properties)sharegroup, + 0 + }; + + cl_int err = 0; + cl::Context cl_context(CL_DEVICE_TYPE_CPU, properties, nullptr, nullptr, &err); + cl::checky(err); + + return cl_context; +} + +cl::Program App::create_cl_program(string file){ + cl_int err = 0; + + // build the program + cl::Program cl_program(cl_context, slurp(file), true, &err); + cout << cl_program.getBuildInfo(cl_device) << std::endl; + + cl::checky(err); + cout << cl_program << endl; + + return cl_program; +} diff --git a/lib/app.hpp b/lib/app.hpp new file mode 100644 index 0000000..6b8227d --- /dev/null +++ b/lib/app.hpp @@ -0,0 +1,51 @@ +#pragma once + +#include +#include +#include +#include + +#include +#include + +struct Vertex{ + moggle::vector3 position; + moggle::vector2 tex_coord; +}; + +struct App { + size_t W = 128; + size_t H = 128; + float time = 0; + float wait = 2; + size_t frame = 0; + + App(GLContext & gl_context); + void draw(); + void resize(size_t w, size_t h); + +private: + GLuint gl_image1; + GLuint gl_image2; + moggle::vbo gl_vbo; + moggle::vao gl_vao; + moggle::shader_program gl_program; + + cl::Context cl_context; + cl::Device cl_device; + cl::CommandQueue cl_queue; + cl::ImageGL cl_image1; + cl::ImageGL cl_image2; + cl::Program cl_program; + KernelOp k_initialize; + KernelOp k_hblur; + KernelOp k_vblur; + KernelOp k_update; + + static moggle::shader_program create_shader_from_files(std::string vertex, std::string fragment); + static moggle::vbo create_vbo(std::vector const & data); + static moggle::vao create_vao(moggle::vbo const & vbo); + static GLuint create_gl_texture(GLsizei width, GLsizei height); + static cl::Context create_shared_cl_context(GLContext & gl_context); + cl::Program create_cl_program(std::string file); +}; diff --git a/lib/utils.hpp b/lib/utils.hpp new file mode 100644 index 0000000..b903d98 --- /dev/null +++ b/lib/utils.hpp @@ -0,0 +1,10 @@ +#pragma once + +#include +#include + +inline 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()}; +} diff --git a/resources/Kernel.cl b/resources/Kernel.cl index e29cc7e..67d0e80 100644 --- a/resources/Kernel.cl +++ b/resources/Kernel.cl @@ -11,6 +11,48 @@ kernel void initialize(size_t width, size_t height, write_only image2d_t output) write_imagef(output, coord, color); } +kernel void hblur(size_t width, size_t height, read_only image2d_t input, write_only image2d_t output){ + int x = get_global_id(0); + int y = get_global_id(1); + + const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE + | CLK_ADDRESS_NONE + | CLK_FILTER_NEAREST; + + float4 c = {0, 0, 0, 0}; + const int size = 7; + float weights[] = {0.25, 0.1, 0.1, 0.1, 0.1, 0.1, 0.25}; + for(int i = 0; i < size; ++i){ + int x2 = (x - size/2 + i + width) % width; + int2 read = {x2, y}; + c += read_imagef(input, sampler, read) * weights[i]; + } + + int2 coord = {x, y}; + write_imagef(output, coord, c); +} + +kernel void vblur(size_t width, size_t height, read_only image2d_t input, write_only image2d_t output){ + int x = get_global_id(0); + int y = get_global_id(1); + + const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE + | CLK_ADDRESS_NONE + | CLK_FILTER_NEAREST; + + float4 c = {0, 0, 0, 0}; + int size = 7; + float weights[] = {0.25, 0.1, 0.1, 0.1, 0.1, 0.1, 0.25}; + for(int i = 0; i < size; ++i){ + int y2 = (y - size/2 + i + height) % height; + int2 read = {x, y2}; + c += read_imagef(input, sampler, read) * weights[i]; + } + + int2 coord = {x, y}; + write_imagef(output, coord, c); +} + kernel void update(size_t width, size_t height, read_only image2d_t input, write_only image2d_t output){ size_t x = get_global_id(0); size_t y = get_global_id(1); diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 4a51e85..d883213 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -3,7 +3,7 @@ set (CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) file(GLOB sources "*.cpp") -set(libs ${JCLWrapper_LIBRARIES} ${JNSAppWrapper_LIBRARIES} ${ImageStreams_LIBRARIES} moggle_xxx) +set(libs common ${JCLWrapper_LIBRARIES} ${JNSAppWrapper_LIBRARIES} ${ImageStreams_LIBRARIES} moggle_xxx) message(STATUS "Linking against the following libraries: ${libs}") foreach(source ${sources}) diff --git a/src/main.cpp b/src/main.cpp index 8dcbe96..a986d64 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,219 +1,23 @@ -#include -#include -#include -#include - -#include -#include -#include -#include +#include "app.hpp" -#include -#include -#include +#include -#include -#include -#include -#include #include using namespace std; -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::vector2 tex_coord; -}; - -const Vertex quad[] = { - // 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}} -}; - -struct App { - 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_device; - unique_ptr cl_queue; - unique_ptr cl_image; - unique_ptr cl_program; - unique_ptr k_initialize; - unique_ptr k_update; - - float time = 0; - GLuint gl_texture = 0; - size_t W = 128; - size_t H = 128; - - float wait = 2; - - enum { - POS_ATTR, - TEX_ATTR - }; - - 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"); - - program.attach(vs); - program.attach(fs); - - program.bind_attribute(POS_ATTR, "position"); - 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(TEX_ATTR); - - v.attribute(POS_ATTR, quad_vbo, &Vertex::position); - v.attribute(TEX_ATTR, quad_vbo, &Vertex::tex_coord); - } - - GLuint create_gl_texture(){ - 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_RGBA32F, W, H, 0, GL_RGBA, GL_UNSIGNED_BYTE, nullptr); - 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; - } - - // shared with opengl - void initialize_cl(GLContext & gl_context){ - 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); - - cl_device = make_unique(cl_context->getInfo().front()); - - cl_queue = make_unique(*cl_context, *cl_device, 0, &err); - cl::checky(err); - } - - void load_cl_kernels(){ - cl_int err = 0; - - // build the program - cl_program = make_unique(*cl_context, slurp("Kernel.cl"), true, &err); - cout << cl_program->getBuildInfo(*cl_device) << std::endl; - - cl::checky(err); - cout << *cl_program << endl; - - // grab the kernels - k_initialize = make_unique(*cl_program, "initialize", &err); - cl::checky(err); - k_update = make_unique(*cl_program, "update", &err); - cl::checky(err); - } - - void initialize(GLContext & gl_context){ - // gl part - load_shader(); - create_vao(); - gl_texture = create_gl_texture(); - - // cl part - initialize_cl(gl_context); - load_cl_kernels(); - - cl_int err = 0; - - // obtain cl memory object - cl_image = make_unique(*cl_context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, gl_texture, &err); - cl::checky(err); - std::cout << *cl_image << std::endl; - - cl::checky((*k_initialize)(*cl_queue, W, H, W, H, *cl_image)); - // no need for finish here, as we will always do a finish in draw(), before gl calls - } - - void draw(){ - auto dt = 1/60.0; - time += dt; - if(wait > 0) wait -= dt; - - if(wait <= 0){ - for(int i = 0; i < 1; ++i){ - cl::checky((*k_update)(*cl_queue, W, H, W, H, *cl_image, *cl_image)); - } - } - - cl::checky(cl_queue->finish()); - - moggle::gl::clear_color(0.65f, 0.65f, 0.65f, 1.0f); - moggle::gl::clear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); - - program.use(); - - 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); - } - - void resize(size_t w, size_t h){ - wait = 2; - - W = w; - H = h; - - gl_texture = create_gl_texture(); - - cl_int err = 0; - cl_image = make_unique(*cl_context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, gl_texture, &err); - cl::checky(err); - std::cout << *cl_image << std::endl; - cl::checky((*k_initialize)(*cl_queue, W, H, W, H, *cl_image)); - } -}; - int main() { - App a; - NSAppWrapper app; - + unique_ptr a; app.create_window({ [&](ContextParameters ctxp){ - a.initialize(ctxp.context); + a = make_unique(ctxp.context); }, [&](ContextParameters){ - a.draw(); + a->draw(); }, [&](ContextParameters, CGFloat w, CGFloat h){ - a.resize(size_t(floor(w)), size_t(floor(h))); + a->resize(size_t(floor(w)), size_t(floor(h))); } });