add compute shader (busted)

1 files changed, 77 insertions, 21 deletions

diff --git a/main.c b/main.c
index cc9ed66..49ca8e2 100644
--- a/main.c
+++ b/main.c
@@ -7,6 +7,10 @@
 
 #define UNUSED(a) ((void) (a))
 
+// string utility macros
+#define SS(v) #v
+#define S(v) SS(v)
+
 // number of time records to keep (must be power of two - 1)
 #define MAX_TIMES 0xFF
 
@@ -14,8 +18,7 @@
 #define EVENT_CODE_TIMER 0
 
 #define NUM_PARTICLES 1000
-#define SS(v) #v
-#define S(v) SS(v)
+#define WORKGROUP_SIZE 1
 
 typedef struct {
   const GLuint type;
@@ -44,19 +47,41 @@ verts[] = {
 };
 
 static const char
+cs_src[] =
+  "#version 320 es\n"
+  "\n"
+  "layout(local_size_x = " S(WORKGROUP_SIZE) ") in;\n"
+  "\n"
+  "layout (std430, binding = 1) buffer particles {\n"
+  "  mediump vec4 positions[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 velocities[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 colors[" S(NUM_PARTICLES) "];\n"
+  "};\n"
+  "\n"
+  // "// layout (location = 0) in mediump float delta;\n"
+  // "\n"
+  "void main() {\n"
+  // "  return;\n"
+  "  mediump uint i = gl_GlobalInvocationID.x;\n"
+  "  positions[i].xy += 0.016 * velocities[i].xy;\n"
+  "}\n";
+
+static const char
 vs_src[] =
   "#version 320 es\n"
   "\n"
   "layout (std430, binding = 1) buffer particles {\n"
-  "  highp vec4 positions[" S(NUM_PARTICLES) "];\n"
-  "  highp vec4 velocities[" S(NUM_PARTICLES) "];\n"
-  "  highp vec4 colors[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 positions[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 velocities[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 colors[" S(NUM_PARTICLES) "];\n"
   "};\n"
   "\n"
   "layout (location = 0) in mediump vec3 pos;\n"
+  "out mediump flat int instance_id;\n"
   "\n"
   "void main() {\n"
   "  gl_Position = positions[gl_InstanceID] + 0.025 * vec4(pos, 1);\n"
+  "  instance_id = gl_InstanceID;\n"
   "}\n";
 
 static const char
@@ -64,20 +89,26 @@ fs_src[] =
   "#version 320 es\n"
   "\n"
   "layout (std430, binding = 1) buffer particles {\n"
-  "  highp vec4 positions[" S(NUM_PARTICLES) "];\n"
-  "  highp vec4 velocities[" S(NUM_PARTICLES) "];\n"
-  "  highp vec4 colors[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 positions[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 velocities[" S(NUM_PARTICLES) "];\n"
+  "  mediump vec4 colors[" S(NUM_PARTICLES) "];\n"
   "};\n"
   "\n"
+  "in mediump flat int instance_id;\n"
   "out mediump vec4 color;\n"
   "\n"
   "uniform mediump float time;\n"
   "\n"
   "void main() {\n"
-  "  color = vec4(0.5f + 0.5 * sin(time * 3.141), 0, 0, 1);\n"
+  "  color = vec4(colors[instance_id].rgb, 1) * (0.5f + 0.5 * sin(time * 3.141));\n"
   "}\n";
 
 static const shader_t
+COMPUTE_SHADERS[] = {
+  { GL_COMPUTE_SHADER, cs_src },
+};
+
+static const shader_t
 RENDER_SHADERS[] = {
   { GL_VERTEX_SHADER,   vs_src },
   { GL_FRAGMENT_SHADER, fs_src },
@@ -340,6 +371,20 @@ static Uint32 timer_cb(
   return interval;
 }
 
+static float
+color_map(uint8_t bits) {
+  switch (bits & 0x3) {
+  case 0x3:
+    return 1.0;
+  case 0x2:
+    return 0.75;
+  case 0x1:
+    return 0.5;
+  default:
+    return 0.25;
+  }
+}
+
 static GLuint
 init_particles(void) {
   // create shader storage buffer
@@ -347,8 +392,8 @@ init_particles(void) {
   glGenBuffers(1, &ssbo);
 
   // generate seed data
-  uint8_t seeds[4 * NUM_PARTICLES];
-  if (!RAND_bytes(seeds, 4 * NUM_PARTICLES)) {
+  uint8_t seeds[5 * NUM_PARTICLES];
+  if (!RAND_bytes(seeds, 5 * NUM_PARTICLES)) {
     SDL_Log("RAND_bytes() failed");
     exit(EXIT_FAILURE);
   }
@@ -357,8 +402,8 @@ init_particles(void) {
   float data[3 * 4 * NUM_PARTICLES];
   #pragma omp parallel for
   for (int i = 0; i < NUM_PARTICLES; i++) {
-    float x = 2.0 * seeds[4 * i + 0] / 0xFF - 1.0,
-          y = 2.0 * seeds[4 * i + 1] / 0xFF - 1.0;
+    float x = 2.0 * seeds[5 * i + 0] / 0xFF - 1.0,
+          y = 2.0 * seeds[5 * i + 1] / 0xFF - 1.0;
 
     // SDL_Log("%f,%f", x, y);
 
@@ -369,13 +414,14 @@ init_particles(void) {
     data[4 * i + 3] = 1.0;
 
     // populate random velocity
-    data[4 * (NUM_PARTICLES + i) + 0] = 2.0 * seeds[4 * i + 2] / 0xFF - 1.0;
-    data[4 * (NUM_PARTICLES + i) + 1] = 2.0 * seeds[4 * i + 3] / 0xFF - 1.0;
+    data[4 * (NUM_PARTICLES + i) + 0] = 2.0 * seeds[5 * i + 2] / 0xFF - 1.0;
+    data[4 * (NUM_PARTICLES + i) + 1] = 2.0 * seeds[5 * i + 3] / 0xFF - 1.0;
 
     // populate colors
-    data[4 * (2 * NUM_PARTICLES + i) + 0] = ((i % 3) == 0) ? 1.0 : 0.0;
-    data[4 * (2 * NUM_PARTICLES + i) + 1] = ((i % 3) == 1) ? 1.0 : 0.0;
-    data[4 * (2 * NUM_PARTICLES + i) + 2] = ((i % 3) == 2) ? 1.0 : 0.0;
+    uint8_t cs = seeds[5 * i + 4];
+    data[4 * (2 * NUM_PARTICLES + i) + 0] = color_map(cs & 0x3);
+    data[4 * (2 * NUM_PARTICLES + i) + 1] = color_map((cs >> 2) & 0x3);
+    data[4 * (2 * NUM_PARTICLES + i) + 2] = color_map((cs >> 4) & 0x3);
     data[4 * (2 * NUM_PARTICLES + i) + 3] = 1.0;
   }
 
@@ -384,8 +430,6 @@ init_particles(void) {
   glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(data), data, GL_STATIC_DRAW);
   glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0);
 
-  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, ssbo);
-
   // return buffer
   return ssbo;
 }
@@ -408,6 +452,7 @@ int main(int argc, char *argv[]) {
 
   // generate shader storage buffer
   GLuint ssbo = init_particles();
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, ssbo);
 
   // generate vertex buffer
   GLuint vbo;
@@ -421,7 +466,11 @@ int main(int argc, char *argv[]) {
   // unbind vao
   glBindVertexArray(0);
 
-  // link program, get uniforms
+  // link compute program, get uniforms
+  GLuint compute_prog = link_program(1, COMPUTE_SHADERS);
+  // GLint u_delta = glGetUniformLocation(compute_prog, "delta");
+
+  // link render program, get uniforms
   GLuint render_prog = link_program(2, RENDER_SHADERS);
   GLint u_time = glGetUniformLocation(render_prog, "time");
 
@@ -433,6 +482,13 @@ int main(int argc, char *argv[]) {
     // save start time
     ctx.times[ctx.times_ofs].draw_start = now;
 
+    // disabled, crashing
+/* 
+ *     glUseProgram(compute_prog);
+ *     glDispatchCompute(NUM_PARTICLES / WORKGROUP_SIZE, 1, 1);
+ *     glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT);
+ */ 
+
     // clear screen
     glClearColor(0, 0, 0, 1);
     glClear(GL_COLOR_BUFFER_BIT);