From 11b1f27f9d27b3f497c20205a77f00272fe8f332 Mon Sep 17 00:00:00 2001
From: David Phillips <david@sighup.nz>
Date: Sun, 12 Feb 2017 00:08:10 +1300
Subject: Add rough, hard-coded mock-up to prove trampoline

This needs the hard-coded constants taken out of it asap to merge
with master. The code is strictly temporary as a measure to prove that
the CL trampoline works and we are actually able to get images from
the device.
---
 fractal-gen.c | 34 ++++++++++++++++++++++++++++++++++
 test.cl       | 24 +++++++++++++++++++++---
 trampoline.c  | 49 ++++++++++++++++++++++++++++++++++++++++++++++++-
 trampoline.h  |  3 +++
 4 files changed, 106 insertions(+), 4 deletions(-)

diff --git a/fractal-gen.c b/fractal-gen.c
index 117e2ca..075858a 100644
--- a/fractal-gen.c
+++ b/fractal-gen.c
@@ -1,4 +1,5 @@
 #include <stdio.h>
+#include <stdlib.h>
 
 #include "trampoline.h"
 
@@ -24,8 +25,41 @@ int main() {
 	}
 	fprintf(stderr, "Compiled.\n");
 
+	fprintf(stderr, "Setting kernel arguments... ");
+	if (tramp_set_kernel_args()) {
+		fprintf(stderr, "Failed.\n");
+		return 1;
+	}
+	fprintf(stderr, "Done.\n");
+
+	fprintf(stderr, "Running kernel... ");
+	if (tramp_run_kernel()) {
+		fprintf(stderr, "Failed.\n");
+		return 1;
+	}
+	fprintf(stderr, "Done.\n");
+
+
+	unsigned int *buffer = calloc(sizeof(unsigned int), 1024*1024);
+	if (!buffer) {
+		perror("malloc");
+		return 1;
+	}
+	fprintf(stderr, "Reading data from device... ");
+	if (tramp_copy_data(&buffer)) {
+		fprintf(stderr, "Failed.\n");
+		return 1;
+	}
+	fprintf(stderr, "Done.\n");
+
 	fprintf(stderr, "Destroying CL trampoline... ");
 	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);
+
 	return 0;
 }
diff --git a/test.cl b/test.cl
index cc587b4..5e081f5 100644
--- a/test.cl
+++ b/test.cl
@@ -1,5 +1,23 @@
-__kernel void test(__global float *a, __global float *b)
+__kernel void fractal_gen(__global unsigned int *buffer)
 {
-	unsigned int i = get_global_id(0);
-	b[i] = a[i] * 2.f;
+	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;
+
+	float2 z = (0.0, 0.0);
+	
+	
+	for (i = 0; i < 254; i++) {
+		// if abs(z) > 2
+		if (z.x*z.x + z.y*z.y > 4)
+			break;
+
+		float oldx = z.x;
+		z.x = z.x*z.x - z.y*z.y + a;
+		z.y = 2*oldx*z.y + b;
+	}
+	buffer[x+1024*y] = i;
 }
diff --git a/trampoline.c b/trampoline.c
index f658d11..fbb68ea 100644
--- a/trampoline.c
+++ b/trampoline.c
@@ -8,8 +8,9 @@ static cl_platform_id platform;
 static cl_context context;
 static cl_device_id* devices;
 static cl_uint device_count;
-unsigned int device_in_use;
+static unsigned int device_in_use;
 static cl_command_queue command_queue;
+static cl_mem device_buffer;
 
 static cl_kernel kernel;
 static cl_program program;
@@ -136,7 +137,53 @@ int tramp_compile_kernel()
 	cl_int ret = 0;
 
 	ret = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+	if (ret != CL_SUCCESS)
+		return 1;
+
+	kernel = clCreateKernel(program, "fractal_gen", &ret);
 
 	/* return non-zero on error */
 	return ret != CL_SUCCESS;
 }
+
+int tramp_set_kernel_args()
+{
+	cl_int ret = 0;
+
+	device_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(unsigned int)*1024*1024, NULL, &ret);
+	if (ret != CL_SUCCESS)
+		return 1;
+
+	ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&device_buffer);
+
+	return ret != CL_SUCCESS;
+}
+
+int tramp_run_kernel()
+{
+	cl_event event;
+	cl_int ret = 0;
+	size_t workgroup_sizes[2];
+	workgroup_sizes[0] = 1024;
+	workgroup_sizes[1] = 1024;
+
+	ret = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, workgroup_sizes, NULL, 0, NULL, &event);
+	if (ret != CL_SUCCESS) {
+		fprintf(stderr, "%d",ret);
+		return ret;
+	}
+
+	clReleaseEvent(event);
+	clFinish(command_queue);
+
+	return ret;
+}
+
+int tramp_copy_data(void **buffer)
+{
+	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);
+	clReleaseEvent(event);
+}
diff --git a/trampoline.h b/trampoline.h
index d5608b4..238980d 100644
--- a/trampoline.h
+++ b/trampoline.h
@@ -3,3 +3,6 @@ void tramp_destroy();
 int tramp_load_kernel(const char *filename);
 char *tramp_get_build_log();
 int tramp_compile_kernel();
+int tramp_set_kernel_args();
+int tramp_run_kernel();
+int tramp_copy_data();
-- 
cgit v1.1