diff --git a/CMakeLists.txt b/CMakeLists.txt index 63499225..f1273c74 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/libs/main/core/pivariant.h b/libs/main/core/pivariant.h index 8f222b63..77313cf4 100644 --- a/libs/main/core/pivariant.h +++ b/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() {return PIVariant::pivBool;} template<> inline PIVariant::Type PIVariant::getType() {return PIVariant::pivChar;} diff --git a/libs/main/io_devices/piserial.cpp b/libs/main/io_devices/piserial.cpp index 07fdded6..203d8e2a 100644 --- a/libs/main/io_devices/piserial.cpp +++ b/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 } diff --git a/libs/main/opencl/piopencl.h b/libs/main/opencl/piopencl.h index 6ec71fa6..5c531054 100644 --- a/libs/main/opencl/piopencl.h +++ b/libs/main/opencl/piopencl.h @@ -31,11 +31,13 @@ public: struct Device; struct Platform; class Context; + class Buffer; class Program; class Kernel; typedef PIVector 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 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 Buffer * createBuffer(PIOpenCL::Direction dir, PIVector & container) { + T def = T(); + return createBuffer(dir, &container, Buffer::cVector , PIByteArray(&def, sizeof(T)), container.size()); + } + template Buffer * createBuffer(PIOpenCL::Direction dir, PIDeque & container) { + T def = T(); + return createBuffer(dir, &container, Buffer::cDeque , PIByteArray(&def, sizeof(T)), container.size()); + } + template Buffer * createBuffer(PIOpenCL::Direction dir, PIVector2D & container) { + T def = T(); + return createBuffer(dir, &container, Buffer::cVector2D, PIByteArray(&def, sizeof(T)), container.size()); + } + template 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 programs_; + PIVector 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() << size);} @@ -168,45 +229,26 @@ public: const PIVector & args() const {return args_;} template bool setArgValue(int index, const T & value) {return setArgValueS(index, PIVariant::fromValue(value));} template bool setArgValue(const PIString & arg, const T & value) {return setArgValue(argIndex(arg), value);} - template bool bindArgValue(int index, PIVector & value, PIOpenCL::Direction dir) { - T def; - return bindArgValueV(index, value.size() * sizeof(T), value.data(), sizeof(def), &def, dir); - } - template bool bindArgValue(const PIString & arg, PIVector & value, PIOpenCL::Direction dir) { - return bindArgValue(argIndex(arg), value, dir); - } - template bool bindArgValue(int index, PIDeque & value, PIOpenCL::Direction dir) { - T def; - return bindArgValueV(index, value.size() * sizeof(T), value.data(), sizeof(def), &def, dir); - } - template bool bindArgValue(const PIString & arg, PIDeque & value, PIOpenCL::Direction dir) { - return bindArgValue(argIndex(arg), value, dir); - } - template bool bindArgValue(int index, PIVector2D & value, PIOpenCL::Direction dir) { - T def; - return bindArgValueV(index, value.size() * sizeof(T), value.data(), sizeof(def), &def, dir); - } - template bool bindArgValue(const PIString & arg, PIVector2D & value, PIOpenCL::Direction dir) { - return bindArgValue(argIndex(arg), value, dir); - } + bool setArgValue(const PIString & arg, const PIVariant & value) {return setArgValueS(argIndex(arg), value);} + bool bindArgValue(int index, Buffer * buffer); + bool bindArgValue(const PIString & arg, Buffer * buffer) {return bindArgValue(argIndex(arg), buffer);} 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 args_; - PIVector buffers_; PIVector dims; PRIVATE_DECLARATION(PIP_OPENCL_EXPORT) }; + static void init(); static const PIVector & platforms(); static const PIVector devices(); diff --git a/libs/opencl/piopencl.cpp b/libs/opencl/piopencl.cpp index e7ab8316..40ea7e19 100644 --- a/libs/opencl/piopencl.cpp +++ b/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 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 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 *)container)->data(); + case cDeque : return ((PIDeque *)container)->data(); + case cVector2D: return ((PIVector2D*)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; - piCout << "execute" << dims; - ret = clEnqueueNDRangeKernel(context_->PRIVATEWB->queue, PRIVATE->kernel, dims.size(), 0, dims.data(), 0, 0, 0, 0); + cl_int 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; - 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); + clSetKernelArg(PRIVATE->kernel, index, sizeof(buffer->PRIVATEWB->buffer), &(buffer->PRIVATEWB->buffer)); return true; } @@ -505,8 +574,6 @@ PIOpenCL::KernelArg::KernelArg() { is_pointer = false; arg_type = Float; dims = 1; - bytes = 0; - buffer = data = 0; }