From c8dc778d855b05a057e8d340b06900ba738811df Mon Sep 17 00:00:00 2001
From: David Phillips <david@sighup.nz>
Date: Sun, 12 Feb 2017 19:53:03 +1300
Subject: Don't hardcode stuff

---
 fractal-gen.c | 24 ++++++++++++++----------
 test.cl       | 13 ++++++++-----
 trampoline.c  | 20 +++++++++++++-------
 3 files changed, 35 insertions(+), 22 deletions(-)

diff --git a/fractal-gen.c b/fractal-gen.c
index 075858a..4401fa1 100644
--- a/fractal-gen.c
+++ b/fractal-gen.c
@@ -3,7 +3,14 @@
 
 #include "trampoline.h"
 
-int main() {
+static unsigned long size;
+static unsigned long iterations;
+
+int main()
+{
+	size = 1024;
+	iterations = 102400;
+
 	fprintf(stderr, "Building CL trampoline... ");
 	if (tramp_init()) {
 		fprintf(stderr, "Failed.\n");
@@ -26,7 +33,7 @@ int main() {
 	fprintf(stderr, "Compiled.\n");
 
 	fprintf(stderr, "Setting kernel arguments... ");
-	if (tramp_set_kernel_args()) {
+	if (tramp_set_kernel_args(size, iterations)) {
 		fprintf(stderr, "Failed.\n");
 		return 1;
 	}
@@ -39,14 +46,13 @@ int main() {
 	}
 	fprintf(stderr, "Done.\n");
 
-
-	unsigned int *buffer = calloc(sizeof(unsigned int), 1024*1024);
+	char *buffer = malloc(size*size);
 	if (!buffer) {
-		perror("malloc");
+		perror("host data buffer malloc");
 		return 1;
 	}
 	fprintf(stderr, "Reading data from device... ");
-	if (tramp_copy_data(&buffer)) {
+	if (tramp_copy_data(&buffer, size*size)) {
 		fprintf(stderr, "Failed.\n");
 		return 1;
 	}
@@ -56,10 +62,8 @@ int main() {
 	tramp_destroy();
 	fprintf(stderr, "Blown to smitherines.\n");
 
-	printf("P5\n1024\n1024\n255\n");
-	unsigned int i;
-	for (i = 0; i < 1024*1024; i++)
-		fputc(buffer[i], stdout);
+	printf("P5\n%d\n%d\n255\n",size,size);
+	fwrite(buffer, size*size, 1, stdout);
 
 	return 0;
 }
diff --git a/test.cl b/test.cl
index 5e081f5..a2ecc67 100644
--- a/test.cl
+++ b/test.cl
@@ -1,16 +1,19 @@
-__kernel void fractal_gen(__global unsigned int *buffer)
+__kernel void fractal_gen(
+	__global unsigned char *buffer,
+	const unsigned int size,
+	const unsigned int iterations)
 {
 	unsigned int x = get_global_id(0);
 	unsigned int y = get_global_id(1);
 	unsigned int i = 0;
 
-	float a = -2.5+(((float)x)/(float)1024)*3.5;
-	float b = -1.75+(((float)y)/(float)1024)*3.5;
+	float a = -2.5+(((float)x)/(float)size)*3.5;
+	float b = -1.75+(((float)y)/(float)size)*3.5;
 
 	float2 z = (0.0, 0.0);
 	
 	
-	for (i = 0; i < 254; i++) {
+	for (i = 0; i < iterations; i++) {
 		// if abs(z) > 2
 		if (z.x*z.x + z.y*z.y > 4)
 			break;
@@ -19,5 +22,5 @@ __kernel void fractal_gen(__global unsigned int *buffer)
 		z.x = z.x*z.x - z.y*z.y + a;
 		z.y = 2*oldx*z.y + b;
 	}
-	buffer[x+1024*y] = i;
+	buffer[x+size*y] = (i*255)/iterations;
 }
diff --git a/trampoline.c b/trampoline.c
index fbb68ea..306e217 100644
--- a/trampoline.c
+++ b/trampoline.c
@@ -146,16 +146,22 @@ int tramp_compile_kernel()
 	return ret != CL_SUCCESS;
 }
 
-int tramp_set_kernel_args()
+int tramp_set_kernel_args(unsigned int size, unsigned int iterations)
 {
 	cl_int ret = 0;
 
-	device_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(unsigned int)*1024*1024, NULL, &ret);
+	device_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 8192*8192, NULL, &ret);
 	if (ret != CL_SUCCESS)
 		return 1;
 
-	ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&device_buffer);
+	ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &device_buffer);
+	fprintf(stderr, "first: %d\n", ret);
+	ret = clSetKernelArg(kernel, 1, sizeof(unsigned int), &size);
+	fprintf(stderr, "first: %d\n", ret);
+	ret = clSetKernelArg(kernel, 2, sizeof(unsigned int), &iterations);
+	fprintf(stderr, "first: %d\n", ret);
 
+	/* FIXME check all clSetKernelArg */
 	return ret != CL_SUCCESS;
 }
 
@@ -164,8 +170,8 @@ int tramp_run_kernel()
 	cl_event event;
 	cl_int ret = 0;
 	size_t workgroup_sizes[2];
-	workgroup_sizes[0] = 1024;
-	workgroup_sizes[1] = 1024;
+	workgroup_sizes[0] = 8192;
+	workgroup_sizes[1] = 8192;
 
 	ret = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, workgroup_sizes, NULL, 0, NULL, &event);
 	if (ret != CL_SUCCESS) {
@@ -179,11 +185,11 @@ int tramp_run_kernel()
 	return ret;
 }
 
-int tramp_copy_data(void **buffer)
+int tramp_copy_data(void **buffer, size_t size)
 {
 	cl_event event;
 	cl_int ret = 0;
 
-	ret = clEnqueueReadBuffer(command_queue, device_buffer, CL_TRUE, 0, sizeof(unsigned int)*1024*1024, *buffer, 0, NULL, &event);
+	ret = clEnqueueReadBuffer(command_queue, device_buffer, CL_TRUE, 0, size, *buffer, 0, NULL, &event);
 	clReleaseEvent(event);
 }
-- 
cgit v1.1