hello-compute: New demoHEADmaster

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],