ACID

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);
 }
 

src/main.cpp

@@ -46,7 +46,12 @@ struct App {
 
 	// 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;
@@ -83,69 +88,80 @@ struct App {
 		v.attribute(TEX_ATTR, quad_vbo, &Vertex::tex_coord);
 	}
 
-	GLuint create_texture(){
+	GLuint create_gl_texture(){
-		std::vector<GLubyte> 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_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){
+	// shared with opengl
-		// Fractal part
+	void initialize_cl(GLContext & gl_context){
-		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_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);
-		auto cl_device = cl_context->getInfo<CL_CONTEXT_DEVICES>().front();
 
-		// obtain cl memory object
+		cl_device = make_unique<cl::Device>(cl_context->getInfo<CL_CONTEXT_DEVICES>().front());
-		cl_image = make_unique<cl::ImageGL>(*cl_context, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, gl_texture, &err);
+
+		cl_queue = make_unique<cl::CommandQueue>(*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 = make_unique<cl::Program>(*cl_context, slurp("Kernel.cl"), true, &err);
-		cl::Program program(*cl_context, {KernelSource.c_str(), KernelSource.size()}, true, &err);
+		cout << cl_program->getBuildInfo<CL_PROGRAM_BUILD_LOG>(*cl_device) << std::endl;
-		cout << 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);
-		cout << program << endl;
+		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;
 
-		// grab the kernel
+		// obtain cl memory object
-		KernelOp kernel(program, "initialize", &err);
+		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;
 
-		// create a queue (for the first device)
+		cl::checky((*k_initialize)(*cl_queue, W, H, W, H, *cl_image));
-		cl::CommandQueue queue(*cl_context, cl_device);
+		// no need for finish here, as we will always do a finish in draw(), before gl calls
-		cl::checky(kernel(queue, W, H, W, H, *cl_image));
-		cl::checky(queue.flush());
 	}
 
 	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<CL_CONTEXT_DEVICES>().front());
-
-		// make a lot of data
-		constexpr size_t W = 512;
-		constexpr size_t H = 512;
-		std::vector<cl_float> 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));
-			}
-		}
-
-		// transfer data into buffers
-		cl::Buffer input(context, input_vector.begin(), input_vector.end(), false, true);
-
-		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
+	void resize(size_t w, size_t h){
-		queue.finish();
+		W = w;
+		H = h;
 
-		auto red = input_vector;
+		gl_texture = create_gl_texture();
 
-		// DO IT (in place)
+		cl_int err = 0;
-		for(int i = 0; i < g; ++i){
+		cl_image = make_unique<cl::ImageGL>(*cl_context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, gl_texture, &err);
-			check(kernel(queue, W, H, input, W, input));
+		cl::checky(err);
-		}
+		std::cout << *cl_image << std::endl;
-
+		cl::checky((*k_initialize)(*cl_queue, W, H, W, H, *cl_image));
-		// 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;
 	}
 };
 
-
 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();