diff options
| -rw-r--r-- | Makefile | 2 | ||||
| -rw-r--r-- | main.c | 90 |
2 files changed, 73 insertions, 19 deletions
@@ -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 @@ -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 |