variable radius

4 files changed, 65 insertions, 33 deletions

diff --git a/balls.c b/balls.c
index 85c3be3..7594eb7 100644
--- a/balls.c
+++ b/balls.c
@@ -19,7 +19,8 @@
 #define VERTEX_SHADER "balls.vert"
 #define FRAGMENT_SHADER "balls.frag"
 
-#define RMAX 0.25f /* Maximum radius. */
+#define RMIN 0.05 /* Minimum radius. */
+#define RMAX 0.10 /* Maximum radius. */
 #define VMAX_INIT 0.1 /* Maximum initial velocity. */
 
 enum { WIDTH = 640, HEIGHT = 480 };
@@ -34,6 +35,7 @@ void initGL(int argc, char *argv[]);
 void initCL(void);
 void setPositions(void);
 void setVelocities(void);
+void setRadii(void);
 void configureSharedData(void);
 void setKernelArgs(void);
 void display(void);
@@ -53,7 +55,7 @@ cl_program prog;
 static cl_command_queue queue;
 static cl_kernel moveKernel, collideWallsKernel, genVerticesKernel;
 GLuint vao, vbo;
-cl_mem positions, velocities, vertexBuf;
+cl_mem positions, velocities, radii, vertexBuf;
 
 int
 main(int argc, char *argv[]) {
@@ -63,6 +65,7 @@ main(int argc, char *argv[]) {
 
 	setPositions();
 	setVelocities();
+	setRadii();
 	configureSharedData();
 	setKernelArgs();
 
@@ -214,6 +217,30 @@ setVelocities(void) {
 }
 
 void
+setRadii(void) {
+	float *hostRadii;
+	int i, err;
+
+	/* Generate radii. */
+	if ((hostRadii = malloc(NBALLS*sizeof(float))) == NULL)
+		sysfatal("Failed to allocate radii array.\n");
+	for (i = 0; i < NBALLS; i++)
+		hostRadii[i] = randFloat(RMIN, RMAX);
+
+	/* Create device-side buffer. */
+	radii = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NBALLS*sizeof(float), hostRadii, &err);
+	if (err <0)
+		sysfatal("Failed to allocate radii buffer.\n");
+
+	/* Copy radii to device. */
+	err = clEnqueueWriteBuffer(queue, radii, CL_TRUE, 0, NBALLS*sizeof(float), hostRadii, 0, NULL, NULL);
+	if (err < 0)
+		sysfatal("Failed to copy radii to device.\n");
+
+	free(hostRadii);
+}
+
+void
 configureSharedData(void) {
 	int err;
 
@@ -243,9 +270,11 @@ setKernelArgs(void) {
 
 	err |= clSetKernelArg(collideWallsKernel, 0, sizeof(positions), &positions);
 	err |= clSetKernelArg(collideWallsKernel, 1, sizeof(velocities), &velocities);
+	err |= clSetKernelArg(collideWallsKernel, 2, sizeof(radii), &radii);
 
 	err |= clSetKernelArg(genVerticesKernel, 0, sizeof(positions), &positions);
-	err |= clSetKernelArg(genVerticesKernel, 1, sizeof(vertexBuf), &vertexBuf);
+	err |= clSetKernelArg(genVerticesKernel, 1, sizeof(radii), &radii);
+	err |= clSetKernelArg(genVerticesKernel, 2, sizeof(vertexBuf), &vertexBuf);
 
 	if (err < 0)
 		sysfatal("Failed to set kernel arguments.\n");
@@ -329,6 +358,7 @@ void
 freeCL(void) {
 	clReleaseMemObject(positions);
 	clReleaseMemObject(velocities);
+	clReleaseMemObject(radii);
 	clReleaseMemObject(vertexBuf);
 
 	clReleaseKernel(moveKernel);
diff --git a/balls.cl b/balls.cl
index 9dff007..b48e0ab 100644
--- a/balls.cl
+++ b/balls.cl
@@ -1,5 +1,3 @@
-#define RADIUS 0.15f
-
 float
 min(float a, float b) {
 	if (a < b)
@@ -28,19 +26,22 @@ move(__global float2 *positions, __global float2 *velocities) {
 }
 
 __kernel void
-collideWalls(__global float2 *positions, __global float2 *velocities) {
-	float2 min, max, p, v;
+collideWalls(__global float2 *positions, __global float2 *velocities, __global float *radii) {
 	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;
+	float2 p, v, min, max;
+	float r;
 
 	id = get_global_id(0);
+
 	p = positions[id];
 	v = velocities[id];
+	r = radii[id];
+
+	/* Set bounds. */
+	min.x = -1.0 + r;
+	min.y = -1.0 + r;
+	max.x = 1.0 - r;
+	max.y = 1.0 - r;
 
 	/* Check for collision with bounds. */
 	if (p.x <= min.x || p.x >= max.x) {
@@ -58,19 +59,20 @@ collideWalls(__global float2 *positions, __global float2 *velocities) {
 }
 
 __kernel void
-genVertices(__global float2 *positions, __global float2 *vertices) {
+genVertices(__global float2 *positions, __global float *radii, __global float2 *vertices) {
 	size_t ball, nsegs;
 	float2 center;
-	float theta;
+	float r, theta;
 
 	ball = get_group_id(0);
 	center = positions[ball];
+	r = radii[ball];
 
 	nsegs = get_local_size(0)-2; /* Number of edge segments. */
 	theta = 2.0f * M_PI_F * get_local_id(0) / nsegs;
 
-	vertices[get_global_id(0)].x = center.x + RADIUS * cos(theta);
-	vertices[get_global_id(0)].y = center.y + RADIUS * sin(theta);
+	vertices[get_global_id(0)].x = center.x + r * cos(theta);
+	vertices[get_global_id(0)].y = center.y + r * sin(theta);
 
 	vertices[ball*get_local_size(0)] = center;
 }
diff --git a/balls.h b/balls.h
index db59b50..051c185 100644
--- a/balls.h
+++ b/balls.h
@@ -6,5 +6,7 @@ typedef struct {
 
 int isCollision(float2 p1, float r1, float2 p2, float r2);
 Rectangle insetRect(Rectangle r, float n);
+
+float randFloat(float lo, float hi);
 float2 randPtInRect(Rectangle r);
 float2 randVec(float xmin, float xmax, float ymin, float ymax);
diff --git a/rand.c b/rand.c
index da450f6..f070fd1 100644
--- a/rand.c
+++ b/rand.c
@@ -3,7 +3,20 @@
 
 #include "balls.h"
 
-static float randFloat(float lo, float hi);
+float
+randFloat(float lo, float hi) {
+	float r, diff;
+	static int isInitialized = 0;
+
+	if (!isInitialized) { /* First call. */
+		srand(time(0));
+		isInitialized = 1;
+	}
+
+	r = (float) rand() / RAND_MAX;
+	diff = hi - lo;
+	return lo + r*diff;
+}
 
 float2
 randPtInRect(Rectangle r) {
@@ -19,18 +32,3 @@ randVec(float xmin, float xmax, float ymin, float ymax) {
 	float2 v = {randFloat(xmin, xmax), randFloat(ymin, ymax)};
 	return v;
 }
-
-static float
-randFloat(float lo, float hi) {
-	float r, diff;
-	static int isInitialized = 0;
-
-	if (!isInitialized) { /* First call. */
-		srand(time(0));
-		isInitialized = 1;
-	}
-
-	r = (float) rand() / RAND_MAX;
-	diff = hi - lo;
-	return lo + r*diff;
-}