Merge branch 'master' of https://git.shs.tools/SHS/pip

CMakeLists.txt

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

libs/main/core/pivariant.h

@@ -661,6 +661,7 @@ template<> inline PIVariantTypes::Color PIVariant::value() const {return toColor
 template<> inline PIVariantTypes::IODevice PIVariant::value() const {return toIODevice();}
 template<> inline PIPointd PIVariant::value() const {return toPoint();}
 template<> inline PIRectd PIVariant::value() const {return toRect();}
+template<> inline PIVariant PIVariant::value() const {return *this;}
 
 template<> inline PIVariant PIVariant::fromValue(const bool & v) {return PIVariant(v);}
 template<> inline PIVariant PIVariant::fromValue(const char & v) {return PIVariant(v);}
@@ -691,6 +692,7 @@ template<> inline PIVariant PIVariant::fromValue(const PIPointd & v) {return PIV
 template<> inline PIVariant PIVariant::fromValue(const PIRectd & v) {return PIVariant(v);}
 template<> inline PIVariant PIVariant::fromValue(const PIMathVectord & v) {return PIVariant(v);}
 template<> inline PIVariant PIVariant::fromValue(const PIMathMatrixd & v) {return PIVariant(v);}
+template<> inline PIVariant PIVariant::fromValue(const PIVariant & v) {return PIVariant(v);}
 
 template<> inline PIVariant::Type PIVariant::getType<bool>() {return PIVariant::pivBool;}
 template<> inline PIVariant::Type PIVariant::getType<char>() {return PIVariant::pivChar;}

libs/main/io_devices/piserial.cpp

@@ -709,7 +709,16 @@ int PISerial::readDevice(void * read_to, int max_size) {
 		return PRIVATE->readed;
 #else
 		if (!canRead()) return -1;
-		return ::read(fd, read_to, max_size);
+		int ret = ::read(fd, read_to, max_size);
+		if (ret < 0) {
+			int err = errno;
+			if (err == EBADF || err == EFAULT || err == EINVAL || err == EIO) {
+				PIThread::stop(false);
+				close();
+				return 0;
+			}
+		}
+		return ret;
 #endif
 }
 

libs/main/opencl/piopencl.h

@@ -31,11 +31,13 @@ public:
 	struct Device;
 	struct Platform;
 	class Context;
+	class Buffer;
 	class Program;
 	class Kernel;
 
 	typedef PIVector<Device> DeviceList;
 
+
 	enum AddressQualifier {
 		AddressGlobal,
 		AddressLocal,
@@ -76,6 +78,7 @@ public:
 		Double,
 	};
 
+
 	struct PIP_OPENCL_EXPORT KernelArg {
 		KernelArg();
 		AddressQualifier address_qualifier;
@@ -91,10 +94,9 @@ public:
 	private:
 		friend class Kernel;
 		void init(void * _k, uint index);
-		int bytes;
-		void * buffer, * data;
 	};
 
+
 	struct PIP_OPENCL_EXPORT Device {
 		Device() {id = platform_id = 0; max_compute_units = max_clock_frequency = 0; max_memory_size = 0;}
 		bool isValid() const {return id != 0;}
@@ -110,6 +112,7 @@ public:
 		ullong max_memory_size;
 	};
 
+
 	struct PIP_OPENCL_EXPORT Platform {
 		Platform() {id = 0;}
 		bool isValid() const {return id != 0;}
@@ -123,7 +126,9 @@ public:
 		PIVector<Device> devices;
 	};
 
+
 	class PIP_OPENCL_EXPORT Context {
+		friend class Buffer;
 		friend class Program;
 		friend class Kernel;
 	public:
@@ -132,17 +137,71 @@ public:
 		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);
+		template <typename T> Buffer * createBuffer(PIOpenCL::Direction dir, PIVector  <T> & container) {
+			T def = T();
+			return createBuffer(dir, &container, Buffer::cVector  , PIByteArray(&def, sizeof(T)), container.size());
+		}
+		template <typename T> Buffer * createBuffer(PIOpenCL::Direction dir, PIDeque   <T> & container) {
+			T def = T();
+			return createBuffer(dir, &container, Buffer::cDeque   , PIByteArray(&def, sizeof(T)), container.size());
+		}
+		template <typename T> Buffer * createBuffer(PIOpenCL::Direction dir, PIVector2D<T> & container) {
+			T def = T();
+			return createBuffer(dir, &container, Buffer::cVector2D, PIByteArray(&def, sizeof(T)), container.size());
+		}
+		template <typename T> Buffer * createBuffer(PIOpenCL::Direction dir, uint elements) {
+			T def = T();
+			Buffer * ret = createBuffer(dir, 0, Buffer::cNone, PIByteArray(&def, sizeof(T)), elements);
+			if (ret)
+				ret->clear();
+			return ret;
+		}
 	private:
 		Context();
 		void zero();
 		void deletePrograms();
+		void deleteBuffers();
+		Buffer * createBuffer(PIOpenCL::Direction dir, void * container, int type, PIByteArray def, uint elements);
 		PIVector<Program * > programs_;
+		PIVector<Buffer * > buffers_;
 		PRIVATE_DECLARATION(PIP_OPENCL_EXPORT)
 	};
 
+
+	class PIP_OPENCL_EXPORT Buffer {
+		friend class Context;
+		friend class Kernel;
+	public:
+		~Buffer();
+		bool resize(uint new_elements);
+		void clear();
+		void copyToContainer();
+		void copyFromContainer();
+	private:
+		enum Container {
+			cNone,
+			cVector,
+			cDeque,
+			cVector2D,
+		};
+		Buffer();
+		void zero();
+		bool init();
+		void * containerData();
+		Context * context_;
+		Direction dir;
+		Container type;
+		void * container;
+		PIByteArray def;
+		uint elements;
+		PRIVATE_DECLARATION(PIP_OPENCL_EXPORT)
+	};
+
+
 	class PIP_OPENCL_EXPORT Program {
 		friend class Context;
 		friend class Kernel;
+		friend class Buffer;
 	public:
 		~Program();
 		const PIString & sourceCode() const {return source_;}
@@ -158,8 +217,10 @@ public:
 		PRIVATE_DECLARATION(PIP_OPENCL_EXPORT)
 	};
 
+
 	class PIP_OPENCL_EXPORT Kernel {
 		friend class Program;
+		friend class Buffer;
 	public:
 		bool execute();
 		void setExecuteRange(int size) {setExecuteRanges(PIVector<int>() << size);}
@@ -168,45 +229,26 @@ public:
 		const PIVector<KernelArg> & args() const {return args_;}
 		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) {
+		bool setArgValue(const PIString & arg, const PIVariant & value) {return setArgValueS(argIndex(arg), value);}
-			T def;
+		bool bindArgValue(int index, Buffer * buffer);
-			return bindArgValueV(index, value.size() * sizeof(T), value.data(), sizeof(def), &def, dir);
+		bool bindArgValue(const PIString & arg, Buffer * buffer) {return bindArgValue(argIndex(arg), buffer);}
-		}
-		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 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)
 	};
 
+
 	static void init();
 	static const PIVector<Platform> & platforms();
 	static const PIVector<Device> devices();

libs/opencl/piopencl.cpp

@@ -15,6 +15,11 @@ PRIVATE_DEFINITION_START(PIOpenCL::Context)
 PRIVATE_DEFINITION_END(PIOpenCL::Context)
 
 
+PRIVATE_DEFINITION_START(PIOpenCL::Buffer)
+	cl_mem buffer;
+PRIVATE_DEFINITION_END(PIOpenCL::Buffer)
+
+
 PRIVATE_DEFINITION_START(PIOpenCL::Program)
 	cl_program program;
 PRIVATE_DEFINITION_END(PIOpenCL::Program)
@@ -141,6 +146,7 @@ PIOpenCL::Context::Context() {
 PIOpenCL::Context::~Context() {
 	piCout << "destroy context" << this;
 	deletePrograms();
+	deleteBuffers();
 	if (PRIVATE->queue)
 		clReleaseCommandQueue(PRIVATE->queue);
 	if (PRIVATE->context)
@@ -151,6 +157,7 @@ PIOpenCL::Context::~Context() {
 
 void PIOpenCL::Context::zero() {
 	programs_.clear();
+	buffers_.clear();
 	PRIVATE->context = 0;
 	PRIVATE->queue = 0;
 	PRIVATE->devices.clear();
@@ -158,7 +165,7 @@ void PIOpenCL::Context::zero() {
 
 
 void PIOpenCL::Context::deletePrograms() {
-	piCout << "context: delete" << programs_.size() << "progs";
+	piCout << "context: delete" << programs_.size() << "programs";
 	PIVector<Program * > ptdl = programs_;
 	programs_.clear();
 	piForeach (Program * p, ptdl) {
@@ -167,6 +174,16 @@ void PIOpenCL::Context::deletePrograms() {
 }
 
 
+void PIOpenCL::Context::deleteBuffers() {
+	piCout << "context: delete" << buffers_.size() << "buffers";
+	PIVector<Buffer * > btdl = buffers_;
+	buffers_.clear();
+	piForeach (Buffer * b, btdl) {
+		if (b) delete b;
+	}
+}
+
+
 PIOpenCL::Context * PIOpenCL::Context::create(const PIOpenCL::DeviceList & dl) {
 	if (dl.isEmpty()) return 0;
 	Context * rc = 0;
@@ -275,6 +292,105 @@ PIOpenCL::Program * PIOpenCL::Context::createProgram(const PIString & source, PI
 }
 
 
+PIOpenCL::Buffer * PIOpenCL::Context::createBuffer(PIOpenCL::Direction dir, void * container, int type, PIByteArray def, uint elements) {
+	Buffer * ret = new Buffer();
+	ret->context_ = this;
+	ret->dir = dir;
+	ret->type = (Buffer::Container)type;
+	ret->container = container;
+	ret->def = def;
+	ret->elements = elements;
+	if (!ret->init()) {
+		delete ret;
+		return 0;
+	}
+	buffers_ << ret;
+	return ret;
+}
+
+
+
+
+PIOpenCL::Buffer::Buffer() {
+	zero();
+}
+
+
+PIOpenCL::Buffer::~Buffer() {
+	if (context_)
+		context_->buffers_.removeAll(this);
+	if (PRIVATE->buffer)
+		clReleaseMemObject(PRIVATE->buffer);
+	zero();
+}
+
+
+void PIOpenCL::Buffer::zero() {
+	type = cNone;
+	container = 0;
+	elements = 0;
+	PRIVATE->buffer = 0;
+}
+
+
+bool PIOpenCL::Buffer::init() {
+	cl_int ret = 0;
+	cl_mem_flags f = container ? CL_MEM_COPY_HOST_PTR : 0;
+	switch (dir) {
+	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;
+	}
+	PRIVATE->buffer = clCreateBuffer(context_->PRIVATEWB->context, f, elements * def.size(), container ? containerData() : 0, &ret);
+	if (ret != 0) {
+		piCout << "[PIOpenCL::Buffer]" << "clCreateBuffer error" << ret;
+		return false;
+	}
+	return true;
+}
+
+
+void * PIOpenCL::Buffer::containerData() {
+	if (type == cNone || !container) return 0;
+	switch (type) {
+	case cVector  : return ((PIVector  <uchar>*)container)->data();
+	case cDeque   : return ((PIDeque   <uchar>*)container)->data();
+	case cVector2D: return ((PIVector2D<uchar>*)container)->data();
+	default: break;
+	}
+	return 0;
+}
+
+
+void PIOpenCL::Buffer::clear() {
+	if (!PRIVATE->buffer) return;
+	if (def.isEmpty() || elements == 0) return;
+	cl_int ret = clEnqueueFillBuffer(context_->PRIVATEWB->queue, PRIVATE->buffer, def.data(), def.size_s(), 0, elements * def.size(), 0, 0, 0);
+	if (ret != 0) {
+		piCout << "[PIOpenCL::Buffer]" << "clEnqueueFillBuffer error" << ret;
+	}
+}
+
+
+void PIOpenCL::Buffer::copyToContainer() {
+	if (!PRIVATE->buffer || !container) return;
+	cl_int ret = clEnqueueReadBuffer(context_->PRIVATEWB->queue, PRIVATE->buffer, CL_TRUE, 0, elements * def.size(), containerData(), 0, 0, 0);
+	if (ret != 0) {
+		piCout << "[PIOpenCL::Buffer]" << "clEnqueueReadBuffer error" << ret;
+	}
+}
+
+
+void PIOpenCL::Buffer::copyFromContainer() {
+	if (!PRIVATE->buffer || !container) return;
+	cl_int ret = clEnqueueWriteBuffer(context_->PRIVATEWB->queue, PRIVATE->buffer, CL_TRUE, 0, elements * def.size(), containerData(), 0, 0, 0);
+	if (ret != 0) {
+		piCout << "[PIOpenCL::Buffer]" << "clEnqueueWriteBuffer error" << ret;
+	}
+}
+
+
 
 
 PIOpenCL::Program::Program() {
@@ -326,22 +442,11 @@ bool PIOpenCL::Kernel::execute() {
 		piCout << "[PIOpenCL::Kernel]" << "Error: empty range";
 		return false;
 	}
-	cl_int ret = 0;
+	cl_int ret = clEnqueueNDRangeKernel(context_->PRIVATEWB->queue, PRIVATE->kernel, dims.size(), 0, dims.data(), 0, 0, 0, 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;
 }
 
@@ -359,10 +464,6 @@ 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);
 }
@@ -432,50 +533,18 @@ bool PIOpenCL::Kernel::setArgValueS(int index, const PIVariant & value) {
 }
 
 
-bool PIOpenCL::Kernel::bindArgValueV(int index, uint bytes, void * value, uint def_bytes, void * def_data, Direction dir) {
+bool PIOpenCL::Kernel::bindArgValue(int index, Buffer * buffer) {
+	if (!buffer) return false;
 	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 << "\"";
+		piCout << "[PIOpenCL::Kernel]" << "bindArgValue set buffer to \"" << ka.type_name << ka.arg_name << "\"";
 		return false;
 	}
-	cl_int ret = 0;
+	clSetKernelArg(PRIVATE->kernel, index, sizeof(buffer->PRIVATEWB->buffer), &(buffer->PRIVATEWB->buffer));
-	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;
 }
 
@@ -505,8 +574,6 @@ PIOpenCL::KernelArg::KernelArg() {
 	is_pointer = false;
 	arg_type = Float;
 	dims = 1;
-	bytes = 0;
-	buffer = data = 0;
 }