aboutsummaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorChad Versace <chad@kiwitree.net>2018-09-20 00:49:18 -0700
committerChad Versace <chad@kiwitree.net>2018-09-20 00:49:18 -0700
commit1486b16227f90080c52805006ce2d9984f49cf54 (patch)
tree27f476087284723f7df97b3f5facc5691fcc152b
parent8f2c36cb0baa6ad93956a0c66452ee63de05176a (diff)
downloadopengl-rutabagas-master.zip
opengl-rutabagas-master.tar.xz
hello-compute: New demoHEADmaster
-rw-r--r--hello-compute.c109
-rw-r--r--hello-compute.comp.glsl11
-rw-r--r--meson.build14
3 files changed, 134 insertions, 0 deletions
diff --git a/hello-compute.c b/hello-compute.c
new file mode 100644
index 0000000..a6e81d0
--- /dev/null
+++ b/hello-compute.c
@@ -0,0 +1,109 @@
+// Copyright 2018 Chad Versace <chad@kiwitree.net>
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// * Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include <assert.h>
+#include <stdint.h>
+#include <stdnoreturn.h>
+
+#include <SDL.h>
+#include <SDL_video.h>
+#include <epoxy/gl.h>
+
+#include "util.h"
+
+static const char shader_src[] = {
+#include "hello-compute.comp.h"
+};
+
+static uint32_t window_width;
+static uint32_t window_height;
+
+static const uint32_t image_width = 256;
+static const uint32_t image_height = 256;
+static const uint32_t workgroup_size_x = 4;
+static const uint32_t workgroup_size_y = 4;
+
+static void
+on_resize(ru_engine_t *e, uint32_t w, uint32_t h) {
+ window_width = w;
+ window_height = h;
+}
+
+static void
+draw(ru_engine_t *e) {
+ glBlitFramebuffer(0, 0, image_width, image_height,
+ 0, 0, window_width, window_height,
+ GL_COLOR_BUFFER_BIT, GL_LINEAR);
+}
+
+int
+main(void) {
+ ru_engine_t engine;
+ ru_engine_init(&engine, "hello-compute", 256, 256);
+ engine.draw = draw;
+ engine.draw_mode = RU_DRAW_MODE_ONCE;
+ engine.resize = on_resize;
+
+ GLuint shader = ru_shader_compile_nofail(GL_COMPUTE_SHADER, shader_src, sizeof(shader_src));
+ GLuint prog = ru_program_attach_link_nofail((GLuint[]) { shader, 0 });
+
+ GLuint image = 0;
+ glGenTextures(1, &image);
+
+ glActiveTexture(GL_TEXTURE0);
+ glBindTexture(GL_TEXTURE_2D, image);
+ glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
+ glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
+ glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
+ glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
+ glTexStorage2D(
+ /*target*/ GL_TEXTURE_2D,
+ /*levels*/ 1,
+ /*interalFormat*/ GL_RGBA8,
+ image_width, image_height);
+ glBindImageTexture(
+ /*unit*/ 0,
+ /*texture*/ image,
+ /*level*/ 0,
+ /*layered*/ false,
+ /*layer*/ 0,
+ GL_WRITE_ONLY,
+ GL_RGBA8);
+
+ glUseProgram(prog);
+
+ assert(image_width % workgroup_size_x == 0);
+ assert(image_height % workgroup_size_y == 0);
+ glDispatchCompute(image_width / workgroup_size_x,
+ image_height / workgroup_size_y,
+ /*z*/ 1);
+
+ // Prepare for blit from image to the default framebuffer.
+ GLuint fb = 0;
+ glGenFramebuffers(1, &fb);
+ glBindFramebuffer(GL_READ_FRAMEBUFFER, fb);
+ glFramebufferTexture(GL_READ_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, image, 0);
+
+ return ru_engine_main_loop(&engine);
+}
diff --git a/hello-compute.comp.glsl b/hello-compute.comp.glsl
new file mode 100644
index 0000000..6f3f49c
--- /dev/null
+++ b/hello-compute.comp.glsl
@@ -0,0 +1,11 @@
+#version 310 es
+
+layout(local_size_x=4, local_size_y=4) in;
+
+mediump layout(binding=0, rgba8) writeonly uniform image2D image;
+
+void main() {
+ vec2 intensity = vec2(gl_GlobalInvocationID.xy) / vec2(imageSize(image));
+ vec4 color = vec4(0, intensity.xy, 1);
+ imageStore(image, ivec2(gl_GlobalInvocationID.xy), color);
+}
diff --git a/meson.build b/meson.build
index ce7b542..f4d394a 100644
--- a/meson.build
+++ b/meson.build
@@ -58,6 +58,20 @@ hello_triangle_exe = executable('hello-triangle',
],
)
+hello_compute_exe = executable('hello-compute',
+ 'hello-compute.c',
+ link_with: [util],
+ dependencies: [
+ dep_epoxy,
+ dep_sdl2,
+ declare_dependency(
+ sources: [
+ gen_glsl_c.process('hello-compute.comp.glsl'),
+ ],
+ ),
+ ],
+)
+
fractal_fern_exe = executable('fractal-fern',
'fractal-fern.c',
link_with: [util],