joshua
/
attractor-cl
Archived
1
0
Fork 0
Code Releases Activity
Browse Source

Iterative texture things (lorenz and blur)

master
Joshua Moerman 11 years ago
parent
4c53f46674
commit
e79a0ea7b5
8 changed files with 294 additions and 203 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 1
      CMakeLists.txt
  2. 9
      lib/CMakeLists.txt
  3. 174
      lib/app.cpp
  4. 51
      lib/app.hpp
  5. 10
      lib/utils.hpp
  6. 42
      resources/Kernel.cl
  7. 2
      src/CMakeLists.txt
  8. 208
      src/main.cpp

1
CMakeLists.txt
View File

@ -4,6 +4,7 @@ cmake_minimum_required(VERSION 2.8)
add_definitions(-std=c++1y) add_definitions(-std=c++1y)
add_subdirectory("contrib") add_subdirectory("contrib")
add_subdirectory("lib")
add_subdirectory("src") add_subdirectory("src")
file(GLOB resources "resources/*") file(GLOB resources "resources/*")

9
lib/CMakeLists.txt
View File

@ -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 ".")

174
lib/app.cpp
View File

@ -0,0 +1,174 @@
#include "app.hpp"
#include "utils.hpp"
#include <iostream>
using namespace std;
static const std::vector<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}}
}};
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<CL_CONTEXT_DEVICES>().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<GLint>("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<Vertex> App::create_vbo(const std::vector<Vertex>& data){
moggle::vbo<Vertex> vbo;
vbo.bind(GL_ARRAY_BUFFER);
vbo.data(data);
return vbo;
}
moggle::vao App::create_vao(const moggle::vbo<Vertex>& 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_PROGRAM_BUILD_LOG>(cl_device) << std::endl;
cl::checky(err);
cout << cl_program << endl;
return cl_program;
}

51
lib/app.hpp
View File

@ -0,0 +1,51 @@
#pragma once
#include <moggle/core/gl.hpp>
#include <moggle/core/shader.hpp>
#include <moggle/core/vao.hpp>
#include <moggle/core/vbo.hpp>
#include <cl2.hpp>
#include <NSGLWrapper.hpp>
struct Vertex{
moggle::vector3<GLfloat> position;
moggle::vector2<GLfloat> 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<Vertex> 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<Vertex> create_vbo(std::vector<Vertex> const & data);
static moggle::vao create_vao(moggle::vbo<Vertex> 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);
};

10
lib/utils.hpp
View File

@ -0,0 +1,10 @@
#pragma once
#include <fstream>
#include <string>
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<char>(f), std::istreambuf_iterator<char>()};
}

42
resources/Kernel.cl
View File

@ -11,6 +11,48 @@ kernel void initialize(size_t width, size_t height, write_only image2d_t output)
write_imagef(output, coord, color); 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){ 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 x = get_global_id(0);
size_t y = get_global_id(1); size_t y = get_global_id(1);

2
src/CMakeLists.txt
View File

@ -3,7 +3,7 @@ set (CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
file(GLOB sources "*.cpp") 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}") message(STATUS "Linking against the following libraries: ${libs}")
foreach(source ${sources}) foreach(source ${sources})

208
src/main.cpp
View File

@ -1,219 +1,23 @@
#include <cl2.hpp> #include "app.hpp"
#include <NSWrapper.hpp>
#include <NSGLWrapper.hpp>
#include <png.hpp>
#include <moggle/core/gl.hpp>
#include <moggle/core/shader.hpp>
#include <moggle/core/vao.hpp>
#include <moggle/core/vbo.hpp>
#include <OpenGL/OpenGL.h> #include <NSWrapper.hpp>
#include <OpenGL/CGLDevice.h>
#include <OpenCL/opencl.h>
#include <vector>
#include <iostream>
#include <fstream>
#include <cassert>
#include <memory> #include <memory>
using namespace std; 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<char>(f), std::istreambuf_iterator<char>()};
}
struct Vertex{
moggle::vector3<GLfloat> position;
moggle::vector2<GLfloat> 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<Vertex> quad_vbo;
// ptrs to postpone construction, should clean this up
unique_ptr<cl::Context> cl_context;
unique_ptr<cl::Device> cl_device;
unique_ptr<cl::CommandQueue> cl_queue;
unique_ptr<cl::ImageGL> cl_image;
unique_ptr<cl::Program> cl_program;
unique_ptr<KernelOp> k_initialize;
unique_ptr<KernelOp> 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::Context>(CL_DEVICE_TYPE_CPU, properties, nullptr, nullptr, &err);
cl::checky(err);
cl_device = make_unique<cl::Device>(cl_context->getInfo<CL_CONTEXT_DEVICES>().front());
cl_queue = make_unique<cl::CommandQueue>(*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::Program>(*cl_context, slurp("Kernel.cl"), true, &err);
cout << cl_program->getBuildInfo<CL_PROGRAM_BUILD_LOG>(*cl_device) << std::endl;
cl::checky(err);
cout << *cl_program << endl;
// grab the kernels
k_initialize = make_unique<KernelOp>(*cl_program, "initialize", &err);
cl::checky(err);
k_update = make_unique<KernelOp>(*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::ImageGL>(*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<GLint>("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::ImageGL>(*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() { int main() {
App a;
NSAppWrapper app; NSAppWrapper app;
unique_ptr<App> a;
app.create_window({ app.create_window({
[&](ContextParameters ctxp){ [&](ContextParameters ctxp){
a.initialize(ctxp.context); a = make_unique<App>(ctxp.context);
}, },
[&](ContextParameters){ [&](ContextParameters){
a.draw(); a->draw();
}, },
[&](ContextParameters, CGFloat w, CGFloat h){ [&](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)));
} }
}); });