collideWalls

2 files changed, 99 insertions, 28 deletions

diff --git a/balls.c b/balls.c
index 05925c9..85c3be3 100644
--- a/balls.c
+++ b/balls.c
@@ -14,6 +14,7 @@
 
 #define PROG_FILE "balls.cl"
 #define MOVE_KERNEL_FUNC "move"
+#define COLLIDE_WALLS_KERNEL_FUNC "collideWalls"
 #define GEN_VERTICES_KERNEL_FUNC "genVertices"
 #define VERTEX_SHADER "balls.vert"
 #define FRAGMENT_SHADER "balls.frag"
@@ -35,20 +36,22 @@ void setPositions(void);
 void setVelocities(void);
 void configureSharedData(void);
 void setKernelArgs(void);
-void execKernel(void);
+void display(void);
+void reshape(int w, int h);
+void move(void);
+void collideWalls(void);
+void genVertices(void);
 void freeCL(void);
 void freeGL(void);
 void initShaders(void);
 char *readFile(const char *filename, size_t *size);
 void compileShader(GLint shader);
-void display(void);
-void reshape(int w, int h);
 float2 *noOverlapPositions(int n);
 
 static cl_context context;
 cl_program prog;
 static cl_command_queue queue;
-static cl_kernel moveKernel, genVerticesKernel;
+static cl_kernel moveKernel, collideWallsKernel, genVerticesKernel;
 GLuint vao, vbo;
 cl_mem positions, velocities, vertexBuf;
 
@@ -157,6 +160,9 @@ initCL(void) {
 	genVerticesKernel = clCreateKernel(prog, GEN_VERTICES_KERNEL_FUNC, &err);
 	if (err < 0)
 		sysfatal("Failed to create kernel '%s': %d\n", GEN_VERTICES_KERNEL_FUNC, err);
+	collideWallsKernel = clCreateKernel(prog, COLLIDE_WALLS_KERNEL_FUNC, &err);
+	if (err < 0)
+		sysfatal("Failed to create kernel '%s': %d\n", COLLIDE_WALLS_KERNEL_FUNC, err);
 	moveKernel = clCreateKernel(prog, MOVE_KERNEL_FUNC, &err);
 	if (err < 0)
 		sysfatal("Failed to create kernel '%s': %d\n", MOVE_KERNEL_FUNC, err);
@@ -235,6 +241,9 @@ setKernelArgs(void) {
 	err = clSetKernelArg(moveKernel, 0, sizeof(positions), &positions);
 	err |= clSetKernelArg(moveKernel, 1, sizeof(velocities), &velocities);
 
+	err |= clSetKernelArg(collideWallsKernel, 0, sizeof(positions), &positions);
+	err |= clSetKernelArg(collideWallsKernel, 1, sizeof(velocities), &velocities);
+
 	err |= clSetKernelArg(genVerticesKernel, 0, sizeof(positions), &positions);
 	err |= clSetKernelArg(genVerticesKernel, 1, sizeof(vertexBuf), &vertexBuf);
 
@@ -243,6 +252,31 @@ setKernelArgs(void) {
 }
 
 void
+display(void) {
+	int i;
+
+	move();
+	collideWalls();
+
+	glClear(GL_COLOR_BUFFER_BIT |GL_DEPTH_BUFFER_BIT);
+
+	genVertices();
+
+	glBindVertexArray(vao);
+	for (i = 0; i < NBALLS; i++)
+		glDrawArrays(GL_TRIANGLE_FAN, i*CIRCLE_POINTS, CIRCLE_POINTS);
+
+	glBindVertexArray(0);
+	glutSwapBuffers();
+	glutPostRedisplay();
+}
+
+void
+reshape(int w, int h) {
+	glViewport(0, 0, (GLsizei) w, (GLsizei) h);
+}
+
+void
 move(void) {
 	size_t size;
 	int err;
@@ -254,6 +288,17 @@ move(void) {
 }
 
 void
+collideWalls(void) {
+	size_t size;
+	int err;
+
+	size = NBALLS;
+	err = clEnqueueNDRangeKernel(queue, collideWallsKernel, 1, NULL, &size, NULL, 0, NULL, NULL);
+	if (err < 0)
+		sysfatal("Couldn't enqueue kernel.\n");
+}
+
+void
 genVertices(void) {
 	int err;
 	size_t localSize, globalSize;
@@ -287,6 +332,7 @@ freeCL(void) {
 	clReleaseMemObject(vertexBuf);
 
 	clReleaseKernel(moveKernel);
+	clReleaseKernel(collideWallsKernel);
 	clReleaseKernel(genVerticesKernel);
 
 	clReleaseCommandQueue(queue);
@@ -370,30 +416,6 @@ compileShader(GLint shader) {
 	}
 }
 
-void
-display(void) {
-	int i;
-
-	move();
-
-	glClear(GL_COLOR_BUFFER_BIT |GL_DEPTH_BUFFER_BIT);
-
-	genVertices();
-
-	glBindVertexArray(vao);
-	for (i = 0; i < NBALLS; i++)
-		glDrawArrays(GL_TRIANGLE_FAN, i*CIRCLE_POINTS, CIRCLE_POINTS);
-
-	glBindVertexArray(0);
-	glutSwapBuffers();
-	glutPostRedisplay();
-}
-
-void
-reshape(int w, int h) {
-	glViewport(0, 0, (GLsizei) w, (GLsizei) h);
-}
-
 float2 *
 noOverlapPositions(int n) {
 	float2 *ps;
diff --git a/balls.cl b/balls.cl
index 0c235c1..9dff007 100644
--- a/balls.cl
+++ b/balls.cl
@@ -1,5 +1,24 @@
 #define RADIUS 0.15f
 
+float
+min(float a, float b) {
+	if (a < b)
+		return a;
+	return b;
+}
+
+float
+max(float a, float b) {
+	if (a > b)
+		return a;
+	return b;
+}
+
+float
+clamp(float x, float lo, float hi) {
+	return min(hi, max(x, lo));
+}
+
 __kernel void
 move(__global float2 *positions, __global float2 *velocities) {
 	size_t id;
@@ -9,6 +28,36 @@ move(__global float2 *positions, __global float2 *velocities) {
 }
 
 __kernel void
+collideWalls(__global float2 *positions, __global float2 *velocities) {
+	float2 min, max, p, v;
+	size_t id;
+
+	/* Set bounds. */
+	min.x = -1.0 + RADIUS;
+	min.y = -1.0 + RADIUS;
+	max.x = 1.0 - RADIUS;
+	max.y = 1.0 - RADIUS;
+
+	id = get_global_id(0);
+	p = positions[id];
+	v = velocities[id];
+
+	/* Check for collision with bounds. */
+	if (p.x <= min.x || p.x >= max.x) {
+		p.x = clamp(p.x, min.x, max.x);
+		v.x = -v.x;
+	}
+	if (p.y <= min.y || p.y >= max.y) {
+		p.y = clamp(p.y, min.y, max.y);
+		v.y = -v.y;
+	}
+
+	/* Write back. */
+	positions[id] = p;
+	velocities[id] = v;
+}
+
+__kernel void
 genVertices(__global float2 *positions, __global float2 *vertices) {
 	size_t ball, nsegs;
 	float2 center;