1.1 --- a/src/clray.cc	Mon Jul 12 05:56:47 2010 +0300
     1.2 +++ b/src/clray.cc	Mon Jul 12 07:00:19 2010 +0300
     1.3 @@ -2,6 +2,18 @@
     1.4  #include <assert.h>
     1.5  #include "ocl.h"
     1.6  
     1.7 +struct Sphere {
     1.8 +	cl_float4 pos;
     1.9 +	cl_float radius;
    1.10 +
    1.11 +	cl_float4 color;
    1.12 +} __attribute__((packed));
    1.13 +
    1.14 +struct Ray {
    1.15 +	cl_float4 origin, dir;
    1.16 +} __attribute__((packed));
    1.17 +
    1.18 +
    1.19  int main()
    1.20  {
    1.21  	int data[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
    1.22 @@ -17,10 +29,13 @@
    1.23  	if(!prog.load("test.cl")) {
    1.24  		return 1;
    1.25  	}
    1.26 -	if(!prog.set_arg(0, ARG_RD, sizeof data, data)) {
    1.27 +	if(!prog.set_arg_buffer(0, ARG_WR, sizeof res, res)) {
    1.28  		return 1;
    1.29  	}
    1.30 -	if(!prog.set_arg(1, ARG_WR, sizeof res, res)) {
    1.31 +	if(!prog.set_arg_buffer(1, ARG_RD, sizeof data, data)) {
    1.32 +		return 1;
    1.33 +	}
    1.34 +	if(!prog.set_argi(2, 100)) {
    1.35  		return 1;
    1.36  	}
    1.37  

     2.1 --- a/src/ocl.cc	Mon Jul 12 05:56:47 2010 +0300
     2.2 +++ b/src/ocl.cc	Mon Jul 12 07:00:19 2010 +0300
     2.3 @@ -140,7 +140,7 @@
     2.4  	prog = 0;
     2.5  	kernel = 0;
     2.6  	this->kname = kname;
     2.7 -	mbuf.resize(16);
     2.8 +	args.resize(16);
     2.9  	built = false;
    2.10  }
    2.11  
    2.12 @@ -154,9 +154,9 @@
    2.13  
    2.14  		clReleaseKernel(kernel);
    2.15  	}
    2.16 -	for(size_t i=0; i<mbuf.size(); i++) {
    2.17 -		if(mbuf[i]) {
    2.18 -			destroy_mem_buffer(mbuf[i]);
    2.19 +	for(size_t i=0; i<args.size(); i++) {
    2.20 +		if(args[i].type == ARGTYPE_MEM_BUF) {
    2.21 +			destroy_mem_buffer(args[i].v.mbuf);
    2.22  		}
    2.23  	}
    2.24  }
    2.25 @@ -192,7 +192,31 @@
    2.26  	return true;
    2.27  }
    2.28  
    2.29 -bool CLProgram::set_arg(int arg, int rdwr, size_t sz, void *ptr)
    2.30 +bool CLProgram::set_argi(int idx, int val)
    2.31 +{
    2.32 +	if((int)args.size() <= idx) {
    2.33 +		args.resize(idx + 1);
    2.34 +	}
    2.35 +
    2.36 +	CLArg *arg = &args[idx];
    2.37 +	arg->type = ARGTYPE_INT;
    2.38 +	arg->v.ival = val;
    2.39 +	return true;
    2.40 +}
    2.41 +
    2.42 +bool CLProgram::set_argf(int idx, float val)
    2.43 +{
    2.44 +	if((int)args.size() <= idx) {
    2.45 +		args.resize(idx + 1);
    2.46 +	}
    2.47 +
    2.48 +	CLArg *arg = &args[idx];
    2.49 +	arg->type = ARGTYPE_FLOAT;
    2.50 +	arg->v.fval = val;
    2.51 +	return true;
    2.52 +}
    2.53 +
    2.54 +bool CLProgram::set_arg_buffer(int idx, int rdwr, size_t sz, void *ptr)
    2.55  {
    2.56  	CLMemBuffer *buf;
    2.57  
    2.58 @@ -200,19 +224,20 @@
    2.59  		return false;
    2.60  	}
    2.61  
    2.62 -	if((int)mbuf.size() <= arg) {
    2.63 -		mbuf.resize(arg + 1);
    2.64 +	if((int)args.size() <= idx) {
    2.65 +		args.resize(idx + 1);
    2.66  	}
    2.67 -	mbuf[arg] = buf;
    2.68 +	args[idx].type = ARGTYPE_MEM_BUF;
    2.69 +	args[idx].v.mbuf = buf;
    2.70  	return true;
    2.71  }
    2.72  
    2.73  CLMemBuffer *CLProgram::get_arg_buffer(int arg)
    2.74  {
    2.75 -	if(arg < 0 || arg >= (int)mbuf.size()) {
    2.76 +	if(arg < 0 || arg >= (int)args.size() || args[arg].type != ARGTYPE_MEM_BUF) {
    2.77  		return 0;
    2.78  	}
    2.79 -	return mbuf[arg];
    2.80 +	return args[arg].v.mbuf;
    2.81  }
    2.82  
    2.83  bool CLProgram::build()
    2.84 @@ -236,22 +261,53 @@
    2.85  		return false;
    2.86  	}
    2.87  
    2.88 -	for(size_t i=0; i<mbuf.size(); i++) {
    2.89 -		if(!mbuf[i]) break;
    2.90 +	for(size_t i=0; i<args.size(); i++) {
    2.91 +		int err;
    2.92  
    2.93 -		int err;
    2.94 -		if((err = clSetKernelArg(kernel, i, sizeof mbuf[i]->mem, &mbuf[i]->mem)) != 0) {
    2.95 -			fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
    2.96 -			clReleaseProgram(prog);
    2.97 -			clReleaseKernel(kernel);
    2.98 -			prog = 0;
    2.99 -			kernel = 0;
   2.100 -			return false;
   2.101 +		if(args[i].type == ARGTYPE_NONE) {
   2.102 +			break;
   2.103 +		}
   2.104 +
   2.105 +		switch(args[i].type) {
   2.106 +		case ARGTYPE_INT:
   2.107 +			if((err = clSetKernelArg(kernel, i, sizeof(int), &args[i].v.ival)) != 0) {
   2.108 +				fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
   2.109 +				goto fail;
   2.110 +			}
   2.111 +			break;
   2.112 +
   2.113 +		case ARGTYPE_FLOAT:
   2.114 +			if((err = clSetKernelArg(kernel, i, sizeof(float), &args[i].v.fval)) != 0) {
   2.115 +				fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
   2.116 +				goto fail;
   2.117 +			}
   2.118 +			break;
   2.119 +
   2.120 +		case ARGTYPE_MEM_BUF:
   2.121 +			{
   2.122 +				CLMemBuffer *mbuf = args[i].v.mbuf;
   2.123 +
   2.124 +				if((err = clSetKernelArg(kernel, i, sizeof mbuf->mem, &mbuf->mem)) != 0) {
   2.125 +					fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
   2.126 +					goto fail;
   2.127 +				}
   2.128 +			}
   2.129 +			break;
   2.130 +
   2.131 +		default:
   2.132 +			break;
   2.133  		}
   2.134  	}
   2.135  
   2.136  	built = true;
   2.137  	return true;
   2.138 +
   2.139 +fail:
   2.140 +	clReleaseProgram(prog);
   2.141 +	clReleaseKernel(kernel);
   2.142 +	prog = 0;
   2.143 +	kernel = 0;
   2.144 +	return false;
   2.145  }
   2.146  
   2.147  bool CLProgram::run() const

     3.1 --- a/src/ocl.h	Mon Jul 12 05:56:47 2010 +0300
     3.2 +++ b/src/ocl.h	Mon Jul 12 07:00:19 2010 +0300
     3.3 @@ -36,13 +36,32 @@
     3.4  bool write_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *src);
     3.5  bool read_mem_buffer(CLMemBuffer *mbuf, size_t sz, void *dest);
     3.6  
     3.7 +enum {
     3.8 +	ARGTYPE_NONE,
     3.9 +
    3.10 +	ARGTYPE_INT,
    3.11 +	ARGTYPE_FLOAT,
    3.12 +	ARGTYPE_FLOAT4,
    3.13 +	ARGTYPE_MEM_BUF
    3.14 +};
    3.15 +
    3.16 +struct CLArg {
    3.17 +	int type;
    3.18 +	union {
    3.19 +		int ival;
    3.20 +		float fval;
    3.21 +		cl_float4 vval;
    3.22 +		CLMemBuffer *mbuf;
    3.23 +	} v;
    3.24 +};
    3.25 +
    3.26  
    3.27  class CLProgram {
    3.28  private:
    3.29  	std::string kname;
    3.30  	cl_program prog;
    3.31  	cl_kernel kernel;
    3.32 -	std::vector<CLMemBuffer*> mbuf;
    3.33 +	std::vector<CLArg> args;
    3.34  	bool built;
    3.35  
    3.36  public:
    3.37 @@ -51,7 +70,9 @@
    3.38  
    3.39  	bool load(const char *fname);
    3.40  
    3.41 -	bool set_arg(int arg, int rdwr, size_t sz, void *buf);
    3.42 +	bool set_argi(int arg, int val);
    3.43 +	bool set_argf(int arg, float val);
    3.44 +	bool set_arg_buffer(int arg, int rdwr, size_t sz, void *buf);
    3.45  	CLMemBuffer *get_arg_buffer(int arg);
    3.46  
    3.47  	bool build();

     4.1 --- a/test.cl	Mon Jul 12 05:56:47 2010 +0300
     4.2 +++ b/test.cl	Mon Jul 12 07:00:19 2010 +0300
     4.3 @@ -1,6 +1,6 @@
     4.4 -__kernel void test(__global const int *src, __global int *dst)
     4.5 +__kernel void test(__global int *dst, __global const int *src, __global const int foo)
     4.6  {
     4.7  	int idx = get_global_id(0);
     4.8  
     4.9 -	dst[idx] = src[idx] * 2.0;
    4.10 +	dst[idx] = src[idx] * foo;
    4.11  }