summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--main.c98
1 files 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;
@@ -44,19 +47,41 @@ verts[] = {
};
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,20 +89,26 @@ 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 },
{ GL_FRAGMENT_SHADER, fs_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);