From f5793ea10055c2241d48883e4a939d06167f8427 Mon Sep 17 00:00:00 2001 From: Paul Duncan Date: Fri, 12 May 2017 19:22:12 -0400 Subject: add compute shader (busted) --- main.c | 98 ++++++++++++++++++++++++++++++++++++++++++++++++++++-------------- 1 file 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; @@ -43,20 +46,42 @@ verts[] = { 0.0, -1, 0, }; +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,19 +89,25 @@ 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 }, @@ -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); -- cgit v1.2.3