summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--Makefile2
-rw-r--r--main.c90
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