pip_code_model macro now automatic add CMAKE_CURRENT_SOURCE_DIR to relative pathes, no ABSOLUTE need

PIOpenCL first working version

CMakeLists.txt

@@ -3,7 +3,7 @@ cmake_policy(SET CMP0017 NEW) # need include() with .cmake
 project(pip)
 set(pip_MAJOR    2)
 set(pip_MINOR    14)
-set(pip_REVISION 1)
+set(pip_REVISION 2)
 set(pip_SUFFIX   )
 set(pip_COMPANY  SHS)
 set(pip_DOMAIN   org.SHS)

cmake/PIPMacros.cmake

@@ -1,11 +1,11 @@
 #[[
 
- pip_code_model(<out_var> file0 [file1 ...] [OPTIONS opt0 [opt1 ...] ] [ABSOLUTE])
+ pip_code_model(<out_var> file0 [file1 ...] [OPTIONS opt0 [opt1 ...] ])
 
  Generate code model files for source files file0 [file1 ...]
 
  Options you can see by exec "pip_cmg -h"
- If not ABSOLUTE source files will be prepended by CMAKE_CURRENT_SOURCE_DIR
+ Relative files pathes read from CMAKE_CURRENT_SOURCE_DIR
  You should add ${<out_var>} to your target
 
 
@@ -53,15 +53,13 @@ macro(pip_code_model RESULT)
 	set(CCM_OUT ${CMAKE_CURRENT_BINARY_DIR}/ccm_${PROJECT_NAME}.cpp)
 	set(${RESULT} ${${RESULT}} ${CCM_OUT})
 	set(CCM_FILES)
-	if (ABS)
+	foreach(csrc ${CCM_SRC})
-		foreach(csrc ${CCM_SRC})
+		if (IS_ABSOLUTE "${csrc}")
 			list(APPEND CCM_FILES "${csrc}")
-		endforeach()
+		else()
-	else()
-		foreach(csrc ${CCM_SRC})
 			list(APPEND CCM_FILES "${CMAKE_CURRENT_SOURCE_DIR}/${csrc}")
-		endforeach()
+		endif()
-	endif()
+	endforeach()
 	#message(STATUS "CCM = ${RESULT}")
 	if(NOT DEFINED PIP_DLL_DIR)
 		set(PIP_DLL_DIR ${CMAKE_CURRENT_BINARY_DIR})

libs/main/opencl/piopencl.h

@@ -50,6 +50,12 @@ public:
 		AccessNone,
 	};
 
+	enum Direction {
+		Input  = 0x01,
+		Output = 0x02,
+		InputOutput = Input | Output,
+	};
+
 	enum TypeQualifier {
 		TypeConst,
 		TypeRestrict,
@@ -74,6 +80,7 @@ public:
 		KernelArg();
 		AddressQualifier address_qualifier;
 		AccessQualifier access_qualifier;
+		Direction direction;
 		TypeQualifier type_qualifier;
 		PIString arg_name;
 		PIString type_name;
@@ -84,6 +91,8 @@ public:
 	private:
 		friend class Kernel;
 		void init(void * _k, uint index);
+		int bytes;
+		void * buffer, * data;
 	};
 
 	struct PIP_OPENCL_EXPORT Device {
@@ -116,10 +125,12 @@ public:
 
 	class PIP_OPENCL_EXPORT Context {
 		friend class Program;
+		friend class Kernel;
 	public:
 		~Context();
 		static Context * create(const DeviceList & dl);
 		static Context * create(const Device & d) {return create(DeviceList() << d);}
+		static Context * create(const PIString & part_name);
 		Program * createProgram(const PIString & source, PIString * error = 0);
 	private:
 		Context();
@@ -135,7 +146,7 @@ public:
 	public:
 		~Program();
 		const PIString & sourceCode() const {return source_;}
-		const Kernel * kernel(int index = 0) const {return kernels_[index];}
+		Kernel * kernel(int index = 0) const {return kernels_[index];}
 		const PIVector<Kernel * > & kernels() const {return kernels_;}
 	private:
 		Program();
@@ -150,22 +161,49 @@ public:
 	class PIP_OPENCL_EXPORT Kernel {
 		friend class Program;
 	public:
+		bool execute();
+		void setExecuteRange(int size) {setExecuteRanges(PIVector<int>() << size);}
+		void setExecuteRanges(const PIVector<int> & ranges);
 		const PIString & name() const {return name_;}
 		const PIVector<KernelArg> & args() const {return args_;}
-		template <typename T> bool setArgValue(int index, const T & value) {return setArgValueV(index, PIVariant::fromValue(value));}
+		template <typename T> bool setArgValue(int index, const T & value) {return setArgValueS(index, PIVariant::fromValue(value));}
 		template <typename T> bool setArgValue(const PIString & arg, const T & value) {return setArgValue(argIndex(arg), value);}
+		template <typename T> bool bindArgValue(int index, PIVector<T> & value, PIOpenCL::Direction dir) {
+			T def;
+			return bindArgValueV(index, value.size() * sizeof(T), value.data(), sizeof(def), &def, dir);
+		}
+		template <typename T> bool bindArgValue(const PIString & arg, PIVector<T> & value, PIOpenCL::Direction dir) {
+			return bindArgValue(argIndex(arg), value, dir);
+		}
+		template <typename T> bool bindArgValue(int index, PIDeque<T> & value, PIOpenCL::Direction dir) {
+			T def;
+			return bindArgValueV(index, value.size() * sizeof(T), value.data(), sizeof(def), &def, dir);
+		}
+		template <typename T> bool bindArgValue(const PIString & arg, PIDeque<T> & value, PIOpenCL::Direction dir) {
+			return bindArgValue(argIndex(arg), value, dir);
+		}
+		template <typename T> bool bindArgValue(int index, PIVector2D<T> & value, PIOpenCL::Direction dir) {
+			T def;
+			return bindArgValueV(index, value.size() * sizeof(T), value.data(), sizeof(def), &def, dir);
+		}
+		template <typename T> bool bindArgValue(const PIString & arg, PIVector2D<T> & value, PIOpenCL::Direction dir) {
+			return bindArgValue(argIndex(arg), value, dir);
+		}
 	private:
 		Kernel();
 		~Kernel();
 		void zero();
 		bool init();
-		bool setArgValueV(int index, const PIVariant & value);
+		bool setArgValueS(int index, const PIVariant & value);
+		bool bindArgValueV(int index, uint bytes, void * value, uint def_bytes, void * def_data, PIOpenCL::Direction dir);
 		int argIndex(const PIString & an) const;
 		KernelArg argByName(const PIString & an) const;
 		Context * context_;
 		Program * program_;
 		PIString name_;
 		PIVector<KernelArg> args_;
+		PIVector<void*> buffers_;
+		PIVector<size_t> dims;
 		PRIVATE_DECLARATION(PIP_OPENCL_EXPORT)
 	};
 

libs/opencl/piopencl.cpp

@@ -195,6 +195,17 @@ PIOpenCL::Context * PIOpenCL::Context::create(const PIOpenCL::DeviceList & dl) {
 }
 
 
+PIOpenCL::Context * PIOpenCL::Context::create(const PIString & part_name) {
+	PIString pn = part_name.toLowerCase();
+	PIVector<Device> dl = PIOpenCL::devices();
+	piForeachC (Device & d, dl) {
+		if (d.displayText().toLowerCase().contains(pn))
+			return create(d);
+	}
+	return 0;
+}
+
+
 PIOpenCL::Program * PIOpenCL::Context::createProgram(const PIString & source, PIString * error) {
 	if (error) error->clear();
 	if (source.isEmpty()) {
@@ -310,6 +321,36 @@ bool PIOpenCL::Program::initKernels(PIVector<void * > kerns) {
 
 
 
+bool PIOpenCL::Kernel::execute() {
+	if (dims.isEmpty()) {
+		piCout << "[PIOpenCL::Kernel]" << "Error: empty range";
+		return false;
+	}
+	cl_int ret = 0;
+	piCout << "execute" << dims;
+	ret = clEnqueueNDRangeKernel(context_->PRIVATEWB->queue, PRIVATE->kernel, dims.size(), 0, dims.data(), 0, 0, 0, 0);
+	if (ret != 0) {
+		piCout << "[PIOpenCL::Kernel]" << "clEnqueueNDRangeKernel error" << ret;
+		return false;
+	}
+	piForeachC (KernelArg & ka, args_) {
+		if (ka.direction == Output || ka.direction == InputOutput) {
+			ret = clEnqueueReadBuffer(context_->PRIVATEWB->queue, (cl_mem)(ka.buffer), CL_TRUE, 0, ka.bytes, ka.data, 0, 0, 0);
+			if (ret != 0) {
+				piCout << "[PIOpenCL::Kernel]" << "clEnqueueWriteBuffer" << ka.type_name << ka.arg_name << "error" << ret;
+				return false;
+			}
+		}
+	}
+	return true;
+}
+
+
+void PIOpenCL::Kernel::setExecuteRanges(const PIVector<int> & ranges) {
+	dims = ranges.toType<size_t>();
+}
+
+
 PIOpenCL::Kernel::Kernel() {
 	zero();
 	//piCout << "new Kernel" << this;
@@ -318,6 +359,10 @@ PIOpenCL::Kernel::Kernel() {
 
 PIOpenCL::Kernel::~Kernel() {
 	//piCout << "del Kernel" << this;
+	piForeachC (void * b, buffers_) {
+		clReleaseMemObject((cl_mem)b);
+	}
+	buffers_.clear();
 	if (PRIVATE->kernel)
 		clReleaseKernel(PRIVATE->kernel);
 }
@@ -349,17 +394,88 @@ bool PIOpenCL::Kernel::init() {
 		ka.init(PRIVATE->kernel, i);
 		args_ << ka;
 	}
-	piCout << "kname" << kname << na;
+	//piCout << "kname" << kname << na;
 	return true;
 }
 
 
-bool PIOpenCL::Kernel::setArgValueV(int index, const PIVariant & value) {
+template <typename T>
+void setArgV(cl_kernel k, int index, T v) {
+	//piCout << "setArgV" << k << index <<sizeof(v)<< v;
+	clSetKernelArg(k, index, sizeof(v), &v);
+}
+
+bool PIOpenCL::Kernel::setArgValueS(int index, const PIVariant & value) {
 	if (index < 0 || index >= args_.size_s()) {
 		piCout << "[PIOpenCL::Kernel]" << "setArgValue invalid index" << index;
 		return false;
 	}
+	KernelArg & ka(args_[index]);
+	if (ka.dims > 0) {
+		piCout << "[PIOpenCL::Kernel]" << "setArgValue set scalar to \"" << ka.type_name << ka.arg_name << "\"";
+		return false;
+	}
+	switch (ka.arg_type) {
+	case Char  : setArgV(PRIVATE->kernel, index, (cl_char)(value.toInt())); break;
+	case UChar : setArgV(PRIVATE->kernel, index, (cl_uchar)(value.toInt())); break;
+	case Short : setArgV(PRIVATE->kernel, index, (cl_short)(value.toInt())); break;
+	case UShort: setArgV(PRIVATE->kernel, index, (cl_ushort)(value.toInt())); break;
+	case Int   : setArgV(PRIVATE->kernel, index, (cl_int)(value.toInt())); break;
+	case UInt  : setArgV(PRIVATE->kernel, index, (cl_uint)(value.toInt())); break;
+	case Long  : setArgV(PRIVATE->kernel, index, (cl_long)(value.toLLong())); break;
+	case ULong : setArgV(PRIVATE->kernel, index, (cl_ulong)(value.toLLong())); break;
+	case Float : setArgV(PRIVATE->kernel, index, (cl_float)(value.toFloat())); break;
+	case Double: setArgV(PRIVATE->kernel, index, (cl_double)(value.toDouble())); break;
+	default: break;
+	}
+	return true;
+}
 
+
+bool PIOpenCL::Kernel::bindArgValueV(int index, uint bytes, void * value, uint def_bytes, void * def_data, Direction dir) {
+	if (index < 0 || index >= args_.size_s()) {
+		piCout << "[PIOpenCL::Kernel]" << "bindArgValue invalid index" << index;
+		return false;
+	}
+	KernelArg & ka(args_[index]);
+	if (ka.dims <= 0) {
+		piCout << "[PIOpenCL::Kernel]" << "bindArgValue set vector to \"" << ka.type_name << ka.arg_name << "\"";
+		return false;
+	}
+	cl_int ret = 0;
+	cl_mem_flags f = CL_MEM_COPY_HOST_PTR;
+	ka.direction = dir;
+	//piCout << "acc" << ka.type_name << ka.arg_name << ka.access_qualifier << bytes << value;
+	switch (ka.direction) {
+	case Input      : f |= CL_MEM_READ_ONLY ; break;
+	case Output     : f |= CL_MEM_WRITE_ONLY; break;
+	case InputOutput: f |= CL_MEM_READ_WRITE; break;
+	default: break;
+	}
+	cl_mem mem = clCreateBuffer(context_->PRIVATEWB->context, f, bytes, value, &ret);
+	if (ret != 0) {
+		piCout << "[PIOpenCL::Kernel]" << "clCreateBuffer" << ka.type_name << ka.arg_name << "error" << ret;
+		return false;
+	}
+	buffers_ << (void*)mem;
+	ka.bytes = bytes;
+	ka.data  = value;
+	ka.buffer = (void*)mem;
+	if (ka.direction == Input || ka.direction == InputOutput) {
+		ret = clEnqueueWriteBuffer(context_->PRIVATEWB->queue, mem, CL_TRUE, 0, bytes, value, 0, 0, 0);
+		if (ret != 0) {
+			piCout << "[PIOpenCL::Kernel]" << "clEnqueueWriteBuffer" << ka.type_name << ka.arg_name << "error" << ret;
+			return false;
+		}
+	}
+	if (ka.direction == Output || ka.direction == InputOutput) {
+		ret = clEnqueueFillBuffer(context_->PRIVATEWB->queue, mem, def_data, def_bytes, 0, bytes, 0, 0, 0);
+		if (ret != 0) {
+			piCout << "[PIOpenCL::Kernel]" << "clEnqueueFillBuffer" << ka.type_name << ka.arg_name << "error" << ret;
+			return false;
+		}
+	}
+	clSetKernelArg(PRIVATE->kernel, index, sizeof(mem), &mem);
 	return true;
 }
 
@@ -389,6 +505,8 @@ PIOpenCL::KernelArg::KernelArg() {
 	is_pointer = false;
 	arg_type = Float;
 	dims = 1;
+	bytes = 0;
+	buffer = data = 0;
 }
 
 
@@ -441,14 +559,26 @@ void PIOpenCL::KernelArg::init(void * _k, uint index) {
 	case CL_KERNEL_ARG_TYPE_VOLATILE: type_qualifier = TypeVolatile; break;
 	case CL_KERNEL_ARG_TYPE_NONE    : type_qualifier = TypeNone;     break;
 	}
-	base_type_name = type_name;
 	is_pointer = false;
-	if (type_name.endsWith("*")) {
+	base_type_name = type_name;
-		is_pointer = true;
+	base_type_name.removeAll("__global");
-		base_type_name.cutRight(1);
+	dims = piMaxi(0, base_type_name.entries('*') + base_type_name.entries('['));
+	base_type_name.removeAll('*');
+	while (base_type_name.contains('[')) {
+		int i = base_type_name.find('[');
+		base_type_name.remove(i, base_type_name.find(']') - i + 1);
 	}
-	dims = piMaxi(1, base_type_name.right(1).toInt());
+	if (base_type_name ==  "char" ) arg_type =  Char ;
-	if (dims > 1) base_type_name.cutRight(1);
+	if (base_type_name == "uchar" ) arg_type = UChar ;
+	if (base_type_name ==  "short") arg_type =  Short;
+	if (base_type_name == "ushort") arg_type = UShort;
+	if (base_type_name ==  "int"  ) arg_type =  Int  ;
+	if (base_type_name == "uint"  ) arg_type = UInt  ;
+	if (base_type_name ==  "long" ) arg_type =  Long ;
+	if (base_type_name == "ulong" ) arg_type = ULong ;
+	if (base_type_name == "float" ) arg_type = Float ;
+	if (base_type_name == "double") arg_type = Double;
+	//piCout << type_name << base_type_name;
 }