From 5c0ef7173af9cb4a8a3edc1edadf4b6303c3e042 Mon Sep 17 00:00:00 2001 From: Ivan Pelipenko Date: Sun, 15 Nov 2020 11:40:47 +0300 Subject: [PATCH] pip_code_model macro now automatic add CMAKE_CURRENT_SOURCE_DIR to relative pathes, no ABSOLUTE need PIOpenCL first working version --- CMakeLists.txt | 2 +- cmake/PIPMacros.cmake | 16 ++-- libs/main/opencl/piopencl.h | 44 ++++++++++- libs/opencl/piopencl.cpp | 146 ++++++++++++++++++++++++++++++++++-- 4 files changed, 187 insertions(+), 21 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 4ae1c13e..3ebfa3fb 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 14) -set(pip_REVISION 1) +set(pip_REVISION 2) set(pip_SUFFIX ) set(pip_COMPANY SHS) set(pip_DOMAIN org.SHS) diff --git a/cmake/PIPMacros.cmake b/cmake/PIPMacros.cmake index acd97207..32b6822b 100644 --- a/cmake/PIPMacros.cmake +++ b/cmake/PIPMacros.cmake @@ -1,11 +1,11 @@ #[[ - pip_code_model( file0 [file1 ...] [OPTIONS opt0 [opt1 ...] ] [ABSOLUTE]) + pip_code_model( 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 ${} 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() - foreach(csrc ${CCM_SRC}) + else() 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}) diff --git a/libs/main/opencl/piopencl.h b/libs/main/opencl/piopencl.h index 6b9ce2ee..6ec71fa6 100644 --- a/libs/main/opencl/piopencl.h +++ b/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 & 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() << size);} + void setExecuteRanges(const PIVector & ranges); const PIString & name() const {return name_;} const PIVector & args() const {return args_;} - template bool setArgValue(int index, const T & value) {return setArgValueV(index, PIVariant::fromValue(value));} + 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); + } 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 args_; + PIVector buffers_; + PIVector dims; PRIVATE_DECLARATION(PIP_OPENCL_EXPORT) }; diff --git a/libs/opencl/piopencl.cpp b/libs/opencl/piopencl.cpp index 8f7313c5..e7ab8316 100644 --- a/libs/opencl/piopencl.cpp +++ b/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 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 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 & ranges) { + dims = ranges.toType(); +} + + 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 +void setArgV(cl_kernel k, int index, T v) { + //piCout << "setArgV" << k << 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("*")) { - is_pointer = true; - base_type_name.cutRight(1); + base_type_name = type_name; + base_type_name.removeAll("__global"); + 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 (dims > 1) base_type_name.cutRight(1); + if (base_type_name == "char" ) arg_type = Char ; + 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; }