diff --git a/resources/Kernel.cl b/resources/Kernel.cl index e22cd69..87d9887 100644 --- a/resources/Kernel.cl +++ b/resources/Kernel.cl @@ -4,15 +4,24 @@ kernel void initialize(size_t width, size_t height, write_only image2d_t output) size_t y = get_global_id(1); int2 coord = {x, y}; - uint4 color = {255 * x / (width - 1), 255 * y / (height - 1), 0, 255}; + //uint4 color = {255 * x / (width - 1), 255 * y / (height - 1), 0, 255}; + //uint4 color = {255 * (x%2), 120 * (y%2), 0, 255}; + uint4 color = {255, 255, 0, 255}; write_imageui(output, coord, color); } -kernel void square(global float* input, size_t width, global float* 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 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)); + + int2 coord = {x, y}; + + const sampler_t samplerA = CLK_NORMALIZED_COORDS_FALSE + | CLK_ADDRESS_NONE + | CLK_FILTER_NEAREST; + uint4 color = read_imageui(input, samplerA, coord).yzxw; + color.x = 255 - color.x; + write_imageui(output, coord, color); } diff --git a/src/main.cpp b/src/main.cpp index 7aad22b..f9fbafc 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -46,7 +46,12 @@ struct App { // 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; @@ -83,69 +88,80 @@ struct App { 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 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_RGBA, W, H, 0, GL_RGBA, GL_UNSIGNED_BYTE, data.data()); + moggle::gl::texture_image_2d(GL_TEXTURE_2D, 0, GL_RGBA, 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; } - void initialize(GLContext & gl_context){ - // Fractal part - load_shader(); - create_vao(); - - // create a texture for opencl - gl_texture = create_texture(); - - // share the bitch + // 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_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_device = make_unique(cl_context->getInfo().front()); + + cl_queue = make_unique(*cl_context, *cl_device, 0, &err); cl::checky(err); - std::cout << *cl_image << std::endl; // sanity check + } + + void load_cl_kernels(){ + cl_int err = 0; // 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_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); - cout << program << endl; + 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; - // grab the kernel - KernelOp kernel(program, "initialize", &err); + // 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; - // 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()); + 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(){ time += 1/60.0f; + 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); @@ -158,99 +174,36 @@ struct App { moggle::gl::draw_arrays(GL_TRIANGLE_STRIP, 0, 4); } -}; - -struct CLApp { - void initialize(){ - auto context = cl::Context::getDefault(); - cout << context << endl; - cl_int err = 0; - // build the program - auto KernelSource = slurp("Kernel.cl"); - cl::Program program(context, {KernelSource.c_str(), KernelSource.size()}, true, &err); - check(err); - cout << program << endl; - - // grab the kernel - KernelOp kernel(program, "square", &err); - check(err); - - // create a queue - cl::CommandQueue queue(context, context.getInfo().front()); - - // make a lot of data - constexpr size_t W = 512; - constexpr size_t H = 512; - std::vector input_vector(W*H); - - for(int y = 0; y < H; ++y){ - for(int x = 0; x < W; ++x){ - input_vector[x + W*y] = 10 * ((x / double(W) - 0.5) + 1.3371337*(y / double(H) - 0.5)); - } - } + void resize(size_t w, size_t h){ + W = w; + H = h; - // transfer data into buffers - cl::Buffer input(context, input_vector.begin(), input_vector.end(), false, true); + gl_texture = create_gl_texture(); - int r = 10, g = 10, b = 10; - - // DO IT (in place) - for(int i = 0; i < r; ++i){ - check(kernel(queue, W, H, input, W, input)); - } - - // read back - queue.finish(); - - auto red = input_vector; - - // DO IT (in place) - for(int i = 0; i < g; ++i){ - check(kernel(queue, W, H, input, W, input)); - } - - // read back - queue.finish(); - - auto green = input_vector; - - // DO IT (in place) - for(int i = 0; i < b; ++i){ - check(kernel(queue, W, H, input, W, input)); - } - - // read back - queue.finish(); - - auto& blue = input_vector; - - // test - cout << "opencl is done" << endl; - png::ostream<> image(W, H, "test.png"); - for(int i = 0; i < red.size(); ++i){ - image << png::ostream<>::pixel(blue[i], red[i], green[i]); - } - cout << "png is saved" << endl; + 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; - CLApp b; NSAppWrapper app; app.create_window({ [&](ContextParameters ctxp){ - //b.initialize(); a.initialize(ctxp.context); }, [&](ContextParameters){ a.draw(); }, - nullptr + [&](ContextParameters, CGFloat w, CGFloat h){ + a.resize(size_t(floor(w)), size_t(floor(h))); + } }); app.run();