# HG changeset patch
# User John Tsiombikas <nuclear@member.fsf.org>
# Date 1278907219 -10800
# Node ID 0b0e4d18d53f6318cd681b1d11bbc93dc8369e29
# Parent  5767277e049fb59990b2104a124032bfdb5b2ccd
added non-buffer args

diff -r 5767277e049f -r 0b0e4d18d53f src/clray.cc
--- a/src/clray.cc	Mon Jul 12 05:56:47 2010 +0300
+++ b/src/clray.cc	Mon Jul 12 07:00:19 2010 +0300
@@ -2,6 +2,18 @@
 #include <assert.h>
 #include "ocl.h"
 
+struct Sphere {
+	cl_float4 pos;
+	cl_float radius;
+
+	cl_float4 color;
+} __attribute__((packed));
+
+struct Ray {
+	cl_float4 origin, dir;
+} __attribute__((packed));
+
+
 int main()
 {
 	int data[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
@@ -17,10 +29,13 @@
 	if(!prog.load("test.cl")) {
 		return 1;
 	}
-	if(!prog.set_arg(0, ARG_RD, sizeof data, data)) {
+	if(!prog.set_arg_buffer(0, ARG_WR, sizeof res, res)) {
 		return 1;
 	}
-	if(!prog.set_arg(1, ARG_WR, sizeof res, res)) {
+	if(!prog.set_arg_buffer(1, ARG_RD, sizeof data, data)) {
+		return 1;
+	}
+	if(!prog.set_argi(2, 100)) {
 		return 1;
 	}
 
diff -r 5767277e049f -r 0b0e4d18d53f src/ocl.cc
--- a/src/ocl.cc	Mon Jul 12 05:56:47 2010 +0300
+++ b/src/ocl.cc	Mon Jul 12 07:00:19 2010 +0300
@@ -140,7 +140,7 @@
 	prog = 0;
 	kernel = 0;
 	this->kname = kname;
-	mbuf.resize(16);
+	args.resize(16);
 	built = false;
 }
 
@@ -154,9 +154,9 @@
 
 		clReleaseKernel(kernel);
 	}
-	for(size_t i=0; i<mbuf.size(); i++) {
-		if(mbuf[i]) {
-			destroy_mem_buffer(mbuf[i]);
+	for(size_t i=0; i<args.size(); i++) {
+		if(args[i].type == ARGTYPE_MEM_BUF) {
+			destroy_mem_buffer(args[i].v.mbuf);
 		}
 	}
 }
@@ -192,7 +192,31 @@
 	return true;
 }
 
-bool CLProgram::set_arg(int arg, int rdwr, size_t sz, void *ptr)
+bool CLProgram::set_argi(int idx, int val)
+{
+	if((int)args.size() <= idx) {
+		args.resize(idx + 1);
+	}
+
+	CLArg *arg = &args[idx];
+	arg->type = ARGTYPE_INT;
+	arg->v.ival = val;
+	return true;
+}
+
+bool CLProgram::set_argf(int idx, float val)
+{
+	if((int)args.size() <= idx) {
+		args.resize(idx + 1);
+	}
+
+	CLArg *arg = &args[idx];
+	arg->type = ARGTYPE_FLOAT;
+	arg->v.fval = val;
+	return true;
+}
+
+bool CLProgram::set_arg_buffer(int idx, int rdwr, size_t sz, void *ptr)
 {
 	CLMemBuffer *buf;
 
@@ -200,19 +224,20 @@
 		return false;
 	}
 
-	if((int)mbuf.size() <= arg) {
-		mbuf.resize(arg + 1);
+	if((int)args.size() <= idx) {
+		args.resize(idx + 1);
 	}
-	mbuf[arg] = buf;
+	args[idx].type = ARGTYPE_MEM_BUF;
+	args[idx].v.mbuf = buf;
 	return true;
 }
 
 CLMemBuffer *CLProgram::get_arg_buffer(int arg)
 {
-	if(arg < 0 || arg >= (int)mbuf.size()) {
+	if(arg < 0 || arg >= (int)args.size() || args[arg].type != ARGTYPE_MEM_BUF) {
 		return 0;
 	}
-	return mbuf[arg];
+	return args[arg].v.mbuf;
 }
 
 bool CLProgram::build()
@@ -236,22 +261,53 @@
 		return false;
 	}
 
-	for(size_t i=0; i<mbuf.size(); i++) {
-		if(!mbuf[i]) break;
+	for(size_t i=0; i<args.size(); i++) {
+		int err;
 
-		int err;
-		if((err = clSetKernelArg(kernel, i, sizeof mbuf[i]->mem, &mbuf[i]->mem)) != 0) {
-			fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
-			clReleaseProgram(prog);
-			clReleaseKernel(kernel);
-			prog = 0;
-			kernel = 0;
-			return false;
+		if(args[i].type == ARGTYPE_NONE) {
+			break;
+		}
+
+		switch(args[i].type) {
+		case ARGTYPE_INT:
+			if((err = clSetKernelArg(kernel, i, sizeof(int), &args[i].v.ival)) != 0) {
+				fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
+				goto fail;
+			}
+			break;
+
+		case ARGTYPE_FLOAT:
+			if((err = clSetKernelArg(kernel, i, sizeof(float), &args[i].v.fval)) != 0) {
+				fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
+				goto fail;
+			}
+			break;
+
+		case ARGTYPE_MEM_BUF:
+			{
+				CLMemBuffer *mbuf = args[i].v.mbuf;
+
+				if((err = clSetKernelArg(kernel, i, sizeof mbuf->mem, &mbuf->mem)) != 0) {
+					fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
+					goto fail;
+				}
+			}
+			break;
+
+		default:
+			break;
 		}
 	}
 
 	built = true;
 	return true;
+
+fail:
+	clReleaseProgram(prog);
+	clReleaseKernel(kernel);
+	prog = 0;
+	kernel = 0;
+	return false;
 }
 
 bool CLProgram::run() const
diff -r 5767277e049f -r 0b0e4d18d53f src/ocl.h
--- a/src/ocl.h	Mon Jul 12 05:56:47 2010 +0300
+++ b/src/ocl.h	Mon Jul 12 07:00:19 2010 +0300
@@ -36,13 +36,32 @@
 bool write_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *src);
 bool read_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *dest);
 
+enum {
+	ARGTYPE_NONE,
+
+	ARGTYPE_INT,
+	ARGTYPE_FLOAT,
+	ARGTYPE_FLOAT4,
+	ARGTYPE_MEM_BUF
+};
+
+struct CLArg {
+	int type;
+	union {
+		int ival;
+		float fval;
+		cl_float4 vval;
+		CLMemBuffer *mbuf;
+	} v;
+};
+
 
 class CLProgram {
 private:
 	std::string kname;
 	cl_program prog;
 	cl_kernel kernel;
-	std::vector<CLMemBuffer*> mbuf;
+	std::vector<CLArg> args;
 	bool built;
 
 public:
@@ -51,7 +70,9 @@
 
 	bool load(const char *fname);
 
-	bool set_arg(int arg, int rdwr, size_t sz, void *buf);
+	bool set_argi(int arg, int val);
+	bool set_argf(int arg, float val);
+	bool set_arg_buffer(int arg, int rdwr, size_t sz, void *buf);
 	CLMemBuffer *get_arg_buffer(int arg);
 
 	bool build();
diff -r 5767277e049f -r 0b0e4d18d53f test.cl
--- a/test.cl	Mon Jul 12 05:56:47 2010 +0300
+++ b/test.cl	Mon Jul 12 07:00:19 2010 +0300
@@ -1,6 +1,6 @@
-__kernel void test(__global const int *src, __global int *dst)
+__kernel void test(__global int *dst, __global const int *src, __global const int foo)
 {
 	int idx = get_global_id(0);
 
-	dst[idx] = src[idx] * 2.0;
+	dst[idx] = src[idx] * foo;
 }