diff options
author | Paul Duncan <pabs@pablotron.org> | 2017-05-12 20:01:12 -0400 |
---|---|---|
committer | Paul Duncan <pabs@pablotron.org> | 2017-05-12 20:01:12 -0400 |
commit | 64391761785447e91154e13894bef402e2b052f7 (patch) | |
tree | c72aeb96dd7ff7d936267a5048face5f972f1b39 | |
parent | f5793ea10055c2241d48883e4a939d06167f8427 (diff) | |
download | compute-test-64391761785447e91154e13894bef402e2b052f7.tar.bz2 compute-test-64391761785447e91154e13894bef402e2b052f7.zip |
add update_particles() and USE_COMPUTE_SHADER
-rw-r--r-- | main.c | 192 |
1 files changed, 120 insertions, 72 deletions
@@ -17,7 +17,13 @@ // custom SDL_USEREVENT code for timer #define EVENT_CODE_TIMER 0 +// use compute shader to update particles? +#define USE_COMPUTE_SHADER false + +// number of particles #define NUM_PARTICLES 1000 + +// number of compute shader workgroups #define WORKGROUP_SIZE 1 typedef struct { @@ -100,7 +106,8 @@ fs_src[] = "uniform mediump float time;\n" "\n" "void main() {\n" - " color = vec4(colors[instance_id].rgb, 1) * (0.5f + 0.5 * sin(time * 3.141));\n" + " mediump float alpha = 0.5f + 0.5f * (0.5f + 0.5 * sin(time * 3.141));\n" + " color = vec4(colors[instance_id].rgb, 1) * alpha;\n" "}\n"; static const shader_t @@ -183,6 +190,113 @@ link_program( return r; } +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 + GLuint ssbo; + glGenBuffers(1, &ssbo); + + // generate seed data + uint8_t seeds[5 * NUM_PARTICLES]; + if (!RAND_bytes(seeds, 5 * NUM_PARTICLES)) { + SDL_Log("RAND_bytes() failed"); + exit(EXIT_FAILURE); + } + + // populate data + float data[3 * 4 * NUM_PARTICLES]; + #pragma omp parallel for + for (int i = 0; i < NUM_PARTICLES; i++) { + 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); + + // populate position + data[4 * i + 0] = x; + data[4 * i + 1] = y; + data[4 * i + 2] = 0.0; + data[4 * i + 3] = 1.0; + + // populate random velocity + 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 + 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; + } + + // populate buffer + glBindBuffer(GL_SHADER_STORAGE_BUFFER, ssbo); + glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(data), data, GL_DYNAMIC_DRAW); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0); + + // return buffer + return ssbo; +} + +static void +update_particles( + const GLuint ssbo, + const GLuint compute_prog +) { + if (USE_COMPUTE_SHADER) { + UNUSED(ssbo); + glUseProgram(compute_prog); + glDispatchCompute(NUM_PARTICLES / WORKGROUP_SIZE, 1, 1); + glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT); + } else { + UNUSED(compute_prog); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, ssbo); + + // map buffer + float *data = glMapBufferRange( + GL_SHADER_STORAGE_BUFFER, + 0, + 3 * NUM_PARTICLES * 4 * sizeof(float), + GL_MAP_READ_BIT | GL_MAP_WRITE_BIT + ); + + #pragma omp parallel for + for (int i = 0; i < NUM_PARTICLES; i++) { + data[4 * i + 0] += 0.016 * data[4 * (NUM_PARTICLES + i) + 0]; + if (data[4 * i + 0] < -1.1) { + data[4 * i + 0] += 2.2; + } else if (data[4 * i + 0] > 1.1) { + data[4 * i + 0] -= 2.2; + } + + data[4 * i + 1] += 0.016 * data[4 * (NUM_PARTICLES + i) + 1]; + if (data[4 * i + 1] < -1.1) { + data[4 * i + 1] += 2.2; + } else if (data[4 * i + 1] > 1.1) { + data[4 * i + 1] -= 2.2; + } + } + + glUnmapBuffer(GL_SHADER_STORAGE_BUFFER); + glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0); + } +} + static void update_viewport( SDL_Window *win @@ -350,7 +464,8 @@ handle_events( } } -static Uint32 timer_cb( +static Uint32 +timer_cb( Uint32 interval, void *arg ) { @@ -371,69 +486,6 @@ 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 - GLuint ssbo; - glGenBuffers(1, &ssbo); - - // generate seed data - uint8_t seeds[5 * NUM_PARTICLES]; - if (!RAND_bytes(seeds, 5 * NUM_PARTICLES)) { - SDL_Log("RAND_bytes() failed"); - exit(EXIT_FAILURE); - } - - // populate data - float data[3 * 4 * NUM_PARTICLES]; - #pragma omp parallel for - for (int i = 0; i < NUM_PARTICLES; i++) { - 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); - - // populate position - data[4 * i + 0] = x; - data[4 * i + 1] = y; - data[4 * i + 2] = 0.0; - data[4 * i + 3] = 1.0; - - // populate random velocity - 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 - 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; - } - - // populate buffer - glBindBuffer(GL_SHADER_STORAGE_BUFFER, ssbo); - glBufferData(GL_SHADER_STORAGE_BUFFER, sizeof(data), data, GL_STATIC_DRAW); - glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0); - - // return buffer - return ssbo; -} - int main(int argc, char *argv[]) { UNUSED(argc); UNUSED(argv); @@ -466,7 +518,7 @@ int main(int argc, char *argv[]) { // unbind vao glBindVertexArray(0); - // link compute program, get uniforms + // link compute program GLuint compute_prog = link_program(1, COMPUTE_SHADERS); // GLint u_delta = glGetUniformLocation(compute_prog, "delta"); @@ -482,12 +534,8 @@ 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); - */ + // update particles + update_particles(ssbo, compute_prog); // clear screen glClearColor(0, 0, 0, 1); |