Uses opencl to generate a texture

resources/Fractal.fsh

@@ -1,26 +1,11 @@
 #version 330
 
-uniform float rotation;
+uniform sampler2D tex;
 
-in vec2 pos;
+in vec2 tex_coord_int;
-in vec2 start;
 
 out vec4 fragColor;
 
 void main(){
-    
+	fragColor = texture(tex, tex_coord_int);
-    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;
 }

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 tex_coord_int;
-out vec2 start;
 
 void main(){
-    pos = position.xy;
+	tex_coord_int = tex_coord;
-    pos.x -= 0.5;
+	gl_Position = position;
-    
-    start = sqrt(1.0 + 0.01 * rotation) * 0.3 * sin(vec2(0.1, 0.1337) * rotation);
-    
-    gl_Position = position;
 }

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 x = get_global_id(0);
-    size_t y = get_global_id(1);
+	size_t y = get_global_id(1);
-    float i = 2.0 * (input[x + width*y] - 0.5);
+	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));
+	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));
 }
 

src/main.cpp

@@ -1,6 +1,6 @@
 #include <cl2.hpp>
-
 #include <NSWrapper.hpp>
+#include <NSGLWrapper.hpp>
 #include <png.hpp>
 
 #include <moggle/core/gl.hpp>
@@ -8,62 +8,57 @@
 #include <moggle/core/vao.hpp>
 #include <moggle/core/vbo.hpp>
 
-#include <OpenGL/gl3.h>
+#include <OpenGL/OpenGL.h>
+#include <OpenGL/CGLDevice.h>
+#include <OpenCL/opencl.h>
 
 #include <vector>
 #include <iostream>
 #include <fstream>
-#include <sstream>
 #include <cassert>
-#include <cmath>
+#include <memory>
 
 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<char>(f), std::istreambuf_iterator<char>()};
+}
 
 struct Vertex{
 	moggle::vector3<GLfloat> position;
-	moggle::vector3<GLfloat> color;
+	moggle::vector2<GLfloat> tex_coord;
 };
 
 const Vertex quad[] = {
-	// x y z, r g b
+	// x   y  z    u  v
-	{{1, -1, 0}, {1, 0, 0}},
+	{{ 1, -1, 0}, {1, 0}},
-	{{-1, -1, 0}, {0, 1, 0}},
+	{{-1, -1, 0}, {0, 0}},
-	{{1, 1, 0}, {0, 0, 1}},
+	{{ 1,  1, 0}, {1, 1}},
-	{{-1, 1, 0}, {1, 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<Vertex> quad_vbo;
 
+	// ptrs to postpone construction, should clean this up
+	unique_ptr<cl::Context> cl_context;
+	unique_ptr<cl::ImageGL> 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<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_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::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_image = make_unique<cl::ImageGL>(*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_PROGRAM_BUILD_LOG>(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<GLfloat>("rotation").set(time);
+
+		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);
@@ -120,8 +180,8 @@ struct CLApp {
 		cl::CommandQueue queue(context, context.getInfo<CL_CONTEXT_DEVICES>().front());
 
 		// make a lot of data
-		constexpr size_t W = 1280 * 1;
+		constexpr size_t W = 512;
-		constexpr size_t H = 800  * 1;
+		constexpr size_t H = 512;
 		std::vector<cl_float> 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){
+		[&](ContextParameters ctxp){
-			a.initialize();
+			//b.initialize();
-			// b.initialize();
+			a.initialize(ctxp.context);
 		},
 		[&](ContextParameters){
 			a.draw();