cleanups, more tests

2 files changed, 73 insertions, 19 deletions

diff --git a/Makefile b/Makefile
index 4294934..cc3adc0 100644
--- a/Makefile
+++ b/Makefile
@@ -1,4 +1,4 @@
-CFLAGS=-std=c11 -fopenmp -W -Wall -pedantic -O2
+CFLAGS=-std=c11 -fopenmp -W -Wall -pedantic -g
 APP=compute-test
 LIBS=-fopenmp -lSDL2 -lGLEW -lGL -lcrypto
 OBJS=main.o
diff --git a/main.c b/main.c
index 9d34f6d..cc4976d 100644
--- a/main.c
+++ b/main.c
@@ -31,6 +31,7 @@
 #define EVENT_CODE_TIMER 0
 
 // use compute shader to update particles?
+// (if false, map the buffer and populate it on the cpu)
 #define USE_COMPUTE_SHADER false
 
 // enable gl debug
@@ -39,8 +40,8 @@
 // number of particles
 #define NUM_PARTICLES 1024
 
-// number of compute shader workgroups
-#define WORKGROUP_SIZE 1024
+// size of compute shader workgroups
+#define WORKGROUP_SIZE 1
 
 typedef struct {
   const GLuint type;
@@ -77,14 +78,15 @@ typedef struct {
 } particles_t;
 
 // shared particle shader storage buffer definition
-#define PARTICLE_BUFFER_DECLARATION \
+#define PARTICLE_BUFFER \
   "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"
 
-// compute shader
+#if 0
+// real compute shader
 static const char
 cs_src[] =
   "#version 320 es\n"
@@ -94,19 +96,27 @@ cs_src[] =
   "\n"
   "layout(local_size_x = " S(WORKGROUP_SIZE) ") in;\n"
   "\n"
-  PARTICLE_BUFFER_DECLARATION
+  PARTICLE_BUFFER
   "\n"
   "void main() {\n"
   "  mediump uint i = gl_GlobalInvocationID.x;\n"
   "  positions[i] += 0.016 * velocities[i];\n"
   "}\n";
+#endif /* 0 */
+
+// empty compute shader
+static const char
+cs_src[] =
+  "#version 320 es\n"
+  "layout(local_size_x = " S(WORKGROUP_SIZE) ") in;\n"
+  "void main() {}\n";
 
 // vertex shader
 static const char
 vs_src[] =
   "#version 320 es\n"
   "\n"
-  PARTICLE_BUFFER_DECLARATION
+  PARTICLE_BUFFER
   "\n"
   "layout (location = 0) in mediump vec3 pos;\n"
   "out mediump flat int instance_id;\n"
@@ -121,7 +131,7 @@ static const char
 fs_src[] =
   "#version 320 es\n"
   "\n"
-  PARTICLE_BUFFER_DECLARATION
+  PARTICLE_BUFFER
   "\n"
   "in mediump flat int instance_id;\n"
   "out mediump vec4 color;\n"
@@ -173,6 +183,7 @@ compile_shader(
   GLuint r = glCreateShader(type);
 
   // compile shader
+  // LOG_D("src = \"%s\"", src);
   glShaderSource(r, 1, &src, NULL);
   glCompileShader(r);
 
@@ -199,10 +210,14 @@ link_program(
   GLuint r = glCreateProgram();
 
   // compile shaders
-  // (FIXME: check num_shaders < 128)
-  GLuint ids[128];
-  if (num_shaders >= 128) {
+  // (FIXME: check num_shaders < 8)
+  GLuint ids[8];
+  if (num_shaders >= 8) {
+    LOG_C("too many shaders: %lu", num_shaders);
+    exit(EXIT_FAILURE);
   }
+
+  // create and attach each shader
   for (size_t i = 0; i < num_shaders; i++) {
     // compile shader
     ids[i] = compile_shader(shaders[i].type, shaders[i].src);
@@ -302,13 +317,42 @@ init_particles(void) {
 
   // populate buffer
   glBindBuffer(GL_SHADER_STORAGE_BUFFER, ssbo);
-  glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(data), &data, GL_DYNAMIC_DRAW);
+  glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(particles_t), &data, GL_DYNAMIC_DRAW);
   glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0);
 
   // return buffer
   return ssbo;
 }
 
+#define GL(op) do { \
+  op; \
+  /* get last GL error */ \
+  GLenum err = glGetError(); \
+  \
+  if (err != GL_NO_ERROR) { \
+    switch (err) { \
+    case GL_OUT_OF_MEMORY: \
+      LOG_C("GL Error: out of memory"); \
+      break; \
+    case GL_INVALID_FRAMEBUFFER_OPERATION: \
+      LOG_C("GL Error: invalid framebuffer operation"); \
+      break; \
+    case GL_INVALID_ENUM: \
+      LOG_C("GL Error: invalid enum"); \
+      break; \
+    case GL_INVALID_VALUE: \
+      LOG_C("GL Error: invalid value"); \
+      break; \
+    case GL_INVALID_OPERATION: \
+      LOG_C("GL Error: invalid operation"); \
+      break; \
+    default: \
+      LOG_C("GL Error: unknown error: %u", err); \
+    } \
+    exit(EXIT_FAILURE); \
+  } \
+} while (0)
+
 static void
 update_particles(
   const GLuint ssbo,
@@ -320,7 +364,6 @@ update_particles(
     glUseProgram(compute_prog);
     glDispatchCompute(NUM_PARTICLES / WORKGROUP_SIZE, 1, 1);
     glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT);
-    // glMemoryBarrier(GL_ALL_BARRIER_BITS);
   } else {
     UNUSED(compute_prog);
 
@@ -386,7 +429,7 @@ ctx_init(context_t * const ctx) {
   }
 
   #ifdef _OPENMP
-  LOG_I("OpenMP enabled.  max_threads = %d", omp_get_max_threads());
+  LOG_I("OpenMP enabled. max_threads = %d", omp_get_max_threads());
   #endif /* _OPENMP */
 
   // create window
@@ -457,9 +500,9 @@ ctx_init(context_t * const ctx) {
     glDebugMessageCallback(debug_cb, ctx);
   }
 
+  // print workgroup size limits
   {
     int sizes[3];
-
     for (int i = 0; i < 3; i++) {
       glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_SIZE, i, sizes + i);
     }
@@ -468,6 +511,15 @@ ctx_init(context_t * const ctx) {
     LOG_I("workgroup_sizes = [%d, %d, %d]", sizes[0], sizes[1], sizes[2]);
   }
 
+  // print shader storage block size limit
+  {
+    GLint val;
+    glGetIntegerv(GL_MAX_SHADER_STORAGE_BLOCK_SIZE, &val);
+
+    // log size limit
+    LOG_I("GL_MAX_SHADER_STORAGE_BLOCK_SIZE = %d", val);
+  }
+
   // set viewport size
   // (after GLEW init)
   update_viewport(ctx->win);
@@ -593,15 +645,15 @@ int main(int argc, char *argv[]) {
   // init timer (5 second interval)
   SDL_TimerID timer_id = SDL_AddTimer(5000, timer_cb, NULL);
 
+  // generate shader storage buffer
+  GLuint ssbo = init_particles();
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, ssbo);
+
   // generate vertex array
   GLuint vao;
   glGenVertexArrays(1, &vao);
   glBindVertexArray(vao);
 
-  // generate shader storage buffer
-  GLuint ssbo = init_particles();
-  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, ssbo);
-
   // generate vertex buffer
   GLuint vbo;
   glGenBuffers(1, &vbo);
@@ -616,6 +668,8 @@ int main(int argc, char *argv[]) {
 
   // link compute program
   GLuint compute_prog = link_program(1, COMPUTE_SHADERS);
+  // TODO: need to pass delta in shader storage block eventually
+  // (for now it's hard-coded as 16ms in shader)
   // GLint u_delta = glGetUniformLocation(compute_prog, "delta");
 
   // link render program, get uniforms