SHS/pip
4
3
Fork 0
Code Issues 67 Pull Requests 1 Actions Releases Wiki Activity

PIVariant::fromValue<PIVariant> fix

Browse Source 
PIOpenCL new Buffer class
This commit is contained in:
Пелипенко Иван
2020-12-03 21:38:30 +03:00
parent 8bfd466f6b
commit 5de87eeed1
3 changed files with 191 additions and 80 deletions
Download Patch File Download Diff File Expand all files Collapse all files

177
libs/opencl/piopencl.cpp
View File

@@ -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;
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;
}