Clover Git
OpenCL 1.1 software implementation
|
CPU kernel work-group. More...
#include <kernel.h>
Classes | |
struct | Context |
Public Member Functions | |
CPUKernelWorkGroup (CPUKernel *kernel, KernelEvent *event, CPUKernelEvent *cpu_event, const size_t *work_group_index) | |
Constructor. | |
~CPUKernelWorkGroup () | |
void * | callArgs (std::vector< void * > &locals_to_free) |
Build a structure of arguments. | |
bool | run () |
Run the work-group. | |
void | builtinNotFound (const std::string &name) const |
Function called when a built-in name cannot be found. | |
Private Member Functions | |
template<typename T > | |
void | writeImageImpl (Image2D *image, int x, int y, int z, T *color) const |
template<typename T > | |
void | readImageImplI (T *result, Image2D *image, int x, int y, int z, uint32_t sampler) const |
template<typename T > | |
void | readImageImplF (T *result, Image2D *image, float x, float y, float z, uint32_t sampler) const |
template<typename T > | |
void | linear3D (T *result, float a, float b, float c, int i0, int j0, int k0, int i1, int j1, int k1, Image3D *image) const |
template<typename T > | |
void | linear2D (T *result, float a, float b, float c, int i0, int j0, int i1, int j1, Image2D *image) const |
Context * | getContextAddr (unsigned int index) |
Private Attributes | |
CPUKernel * | p_kernel |
CPUKernelEvent * | p_cpu_event |
KernelEvent * | p_event |
cl_uint | p_work_dim |
size_t | p_index [MAX_WORK_DIMS] |
size_t | p_max_local_id [MAX_WORK_DIMS] |
size_t | p_global_id_start_offset [MAX_WORK_DIMS] |
void(* | p_kernel_func_addr )(void *) |
void * | p_args |
Context * | p_current_context |
Context | p_dummy_context |
void * | p_contexts |
size_t | p_stack_size |
unsigned int | p_num_work_items |
unsigned int | p_current_work_item |
bool | p_had_barrier |
Native implementation of built-in OpenCL C functions | |
size_t | getGlobalId (cl_uint dimindx) const |
cl_uint | getWorkDim () const |
size_t | getGlobalSize (cl_uint dimindx) const |
size_t | getLocalSize (cl_uint dimindx) const |
size_t | getLocalID (cl_uint dimindx) const |
size_t | getNumGroups (cl_uint dimindx) const |
size_t | getGroupID (cl_uint dimindx) const |
size_t | getGlobalOffset (cl_uint dimindx) const |
void | barrier (unsigned int flags) |
void * | getImageData (Image2D *image, int x, int y, int z) const |
void | writeImage (Image2D *image, int x, int y, int z, float *color) const |
void | writeImage (Image2D *image, int x, int y, int z, int32_t *color) const |
void | writeImage (Image2D *image, int x, int y, int z, uint32_t *color) const |
void | readImage (float *result, Image2D *image, int x, int y, int z, uint32_t sampler) const |
void | readImage (int32_t *result, Image2D *image, int x, int y, int z, uint32_t sampler) const |
void | readImage (uint32_t *result, Image2D *image, int x, int y, int z, uint32_t sampler) const |
void | readImage (float *result, Image2D *image, float x, float y, float z, uint32_t sampler) const |
void | readImage (int32_t *result, Image2D *image, float x, float y, float z, uint32_t sampler) const |
void | readImage (uint32_t *result, Image2D *image, float x, float y, float z, uint32_t sampler) const |
CPU kernel work-group.
This class represent a bulk of work-items that will be run. It is the one to actually run the kernel of its elements.
Definition at line 150 of file cpu/kernel.h.
CPUKernelWorkGroup::CPUKernelWorkGroup | ( | CPUKernel * | kernel, |
KernelEvent * | event, | ||
CPUKernelEvent * | cpu_event, | ||
const size_t * | work_group_index | ||
) |
Constructor.
kernel | kernel to run |
event | event containing information about the kernel run |
cpu_event | CPU-specific information and cache about event |
work_group_index | index of this work-group in the kernel |
Definition at line 388 of file cpu/kernel.cpp.
References Coal::KernelEvent::global_work_offset(), p_global_id_start_offset, p_index, p_max_local_id, p_num_work_items, and p_work_dim.
CPUKernelWorkGroup::~CPUKernelWorkGroup | ( | ) |
Definition at line 413 of file cpu/kernel.cpp.
References p_cpu_event, and Coal::CPUKernelEvent::workGroupFinished().
void CPUKernelWorkGroup::barrier | ( | unsigned int | flags | ) |
Definition at line 159 of file builtins.cpp.
References Coal::CPUKernelWorkGroup::Context::context, Coal::CPUKernel::function(), getContextAddr(), getWorkItemsData(), incVec(), Coal::CPUKernelWorkGroup::Context::initialized, Coal::CPUKernelWorkGroup::Context::local_id, main(), p_args, p_contexts, p_current_context, p_current_work_item, p_had_barrier, p_kernel, p_kernel_func_addr, p_max_local_id, p_num_work_items, p_stack_size, p_work_dim, and setWorkItemsData().
Referenced by barrier().
void CPUKernelWorkGroup::builtinNotFound | ( | const std::string & | name | ) | const |
Function called when a built-in name cannot be found.
Definition at line 248 of file builtins.cpp.
References Coal::CPUKernel::function(), and p_kernel.
Referenced by getBuiltin().
void * CPUKernelWorkGroup::callArgs | ( | std::vector< void * > & | locals_to_free | ) |
Build a structure of arguments.
As C doesn't support calling functions with variable arguments unknown at the compilation, this function builds the list of arguments in memory. This array will then be passed to a LLVM stub function reading it and passing its values to the actuel kernel.
locals_to_free | if this kernel takes __local arguments, they must be malloc()'ed for every work-group. They are placed in this vector to be free()'ed at the end of run() . |
Definition at line 418 of file cpu/kernel.cpp.
References Coal::MemObject::allocate(), Coal::Kernel::Arg::allocAtKernelRuntime(), Coal::Kernel::arg(), Coal::Kernel::Arg::Buffer, Coal::CPUKernelEvent::cacheKernelArgs(), Coal::CPUBuffer::data(), Coal::Kernel::Arg::data(), Coal::CPUKernel::device(), Coal::MemObject::deviceBuffer(), Coal::Kernel::Arg::file(), Coal::Kernel::hasLocals(), Coal::Kernel::Arg::Image2D, Coal::Kernel::Arg::Image3D, Coal::CPUKernel::kernel(), Coal::CPUKernelEvent::kernelArgs(), Coal::Kernel::Arg::kind(), Coal::Kernel::Arg::Local, Coal::Kernel::numArgs(), p_cpu_event, p_kernel, Coal::CPUKernel::typeOffset(), Coal::Kernel::Arg::valueSize(), and Coal::Kernel::Arg::vecDim().
Referenced by run().
CPUKernelWorkGroup::Context * CPUKernelWorkGroup::getContextAddr | ( | unsigned int | index | ) | [private] |
Definition at line 580 of file cpu/kernel.cpp.
References p_contexts, and p_stack_size.
size_t CPUKernelWorkGroup::getGlobalId | ( | cl_uint | dimindx | ) | const |
Definition at line 102 of file builtins.cpp.
References Coal::CPUKernelWorkGroup::Context::local_id, p_current_context, p_global_id_start_offset, and p_work_dim.
Referenced by get_global_id().
size_t CPUKernelWorkGroup::getGlobalOffset | ( | cl_uint | dimindx | ) | const |
Definition at line 151 of file builtins.cpp.
References Coal::KernelEvent::global_work_offset(), p_event, and p_work_dim.
Referenced by get_global_offset().
size_t CPUKernelWorkGroup::getGlobalSize | ( | cl_uint | dimindx | ) | const |
Definition at line 110 of file builtins.cpp.
References Coal::KernelEvent::global_work_size(), p_event, and p_work_dim.
Referenced by get_global_size().
size_t CPUKernelWorkGroup::getGroupID | ( | cl_uint | dimindx | ) | const |
Definition at line 143 of file builtins.cpp.
References p_index, and p_work_dim.
Referenced by get_group_id().
void * CPUKernelWorkGroup::getImageData | ( | Image2D * | image, |
int | x, | ||
int | y, | ||
int | z | ||
) | const |
Definition at line 421 of file cpu/sampler.cpp.
References Coal::CPUBuffer::data(), Coal::CPUKernel::device(), Coal::MemObject::deviceBuffer(), imageData(), p_kernel, Coal::Image2D::pixel_size(), Coal::Image2D::row_pitch(), and Coal::Image2D::slice_pitch().
Referenced by image_data(), readImageImplI(), and writeImageImpl().
size_t CPUKernelWorkGroup::getLocalID | ( | cl_uint | dimindx | ) | const |
Definition at line 126 of file builtins.cpp.
References Coal::CPUKernelWorkGroup::Context::local_id, p_current_context, and p_work_dim.
Referenced by get_local_id().
size_t CPUKernelWorkGroup::getLocalSize | ( | cl_uint | dimindx | ) | const |
Definition at line 118 of file builtins.cpp.
References Coal::KernelEvent::local_work_size(), p_event, and p_work_dim.
Referenced by get_local_size().
size_t CPUKernelWorkGroup::getNumGroups | ( | cl_uint | dimindx | ) | const |
Definition at line 134 of file builtins.cpp.
References Coal::KernelEvent::global_work_size(), Coal::KernelEvent::local_work_size(), p_event, and p_work_dim.
Referenced by get_num_groups().
cl_uint CPUKernelWorkGroup::getWorkDim | ( | ) | const |
void CPUKernelWorkGroup::linear2D | ( | T * | result, |
float | a, | ||
float | b, | ||
float | c, | ||
int | i0, | ||
int | j0, | ||
int | i1, | ||
int | j1, | ||
Image2D * | image | ||
) | const [private] |
Definition at line 310 of file cpu/sampler.cpp.
References vec4_add(), and vec4_scalar_mul().
void CPUKernelWorkGroup::linear3D | ( | T * | result, |
float | a, | ||
float | b, | ||
float | c, | ||
int | i0, | ||
int | j0, | ||
int | k0, | ||
int | i1, | ||
int | j1, | ||
int | k1, | ||
Image3D * | image | ||
) | const [private] |
Definition at line 271 of file cpu/sampler.cpp.
References vec4_add(), and vec4_scalar_mul().
void CPUKernelWorkGroup::readImage | ( | float * | result, |
Image2D * | image, | ||
int | x, | ||
int | y, | ||
int | z, | ||
uint32_t | sampler | ||
) | const |
Definition at line 537 of file cpu/sampler.cpp.
Referenced by read_imageff(), read_imagefi(), read_imageif(), read_imageii(), read_imageuif(), and read_imageuii().
void CPUKernelWorkGroup::readImage | ( | int32_t * | result, |
Image2D * | image, | ||
int | x, | ||
int | y, | ||
int | z, | ||
uint32_t | sampler | ||
) | const |
Definition at line 543 of file cpu/sampler.cpp.
void CPUKernelWorkGroup::readImage | ( | uint32_t * | result, |
Image2D * | image, | ||
int | x, | ||
int | y, | ||
int | z, | ||
uint32_t | sampler | ||
) | const |
Definition at line 549 of file cpu/sampler.cpp.
void CPUKernelWorkGroup::readImage | ( | int32_t * | result, |
Image2D * | image, | ||
float | x, | ||
float | y, | ||
float | z, | ||
uint32_t | sampler | ||
) | const |
Definition at line 759 of file cpu/sampler.cpp.
void CPUKernelWorkGroup::readImage | ( | float * | result, |
Image2D * | image, | ||
float | x, | ||
float | y, | ||
float | z, | ||
uint32_t | sampler | ||
) | const |
Definition at line 753 of file cpu/sampler.cpp.
void CPUKernelWorkGroup::readImage | ( | uint32_t * | result, |
Image2D * | image, | ||
float | x, | ||
float | y, | ||
float | z, | ||
uint32_t | sampler | ||
) | const |
Definition at line 765 of file cpu/sampler.cpp.
void CPUKernelWorkGroup::readImageImplF | ( | T * | result, |
Image2D * | image, | ||
float | x, | ||
float | y, | ||
float | z, | ||
uint32_t | sampler | ||
) | const [private] |
Definition at line 556 of file cpu/sampler.cpp.
References CLK_ADDRESS_CLAMP, CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_MIRRORED_REPEAT, CLK_ADDRESS_NONE, CLK_ADDRESS_REPEAT, CLK_FILTER_LINEAR, CLK_FILTER_NEAREST, CLK_NORMALIZED_COORDS_TRUE, Coal::Image3D::depth(), frac(), Coal::Image2D::height(), Coal::MemObject::Image3D, max(), min(), round(), Coal::Image2D::type(), and Coal::Image2D::width().
void CPUKernelWorkGroup::readImageImplI | ( | T * | result, |
Image2D * | image, | ||
int | x, | ||
int | y, | ||
int | z, | ||
uint32_t | sampler | ||
) | const [private] |
Definition at line 497 of file cpu/sampler.cpp.
References Coal::Image2D::channels(), convert_from_format(), Coal::Image2D::format(), getImageData(), handle_address_mode(), and swizzle().
bool CPUKernelWorkGroup::run | ( | ) |
Run the work-group.
This function is the core of CPU-acceleration. It runs the work-items of this work-group given the correct arguments.
Definition at line 514 of file cpu/kernel.cpp.
References callArgs(), Coal::CPUKernel::callFunction(), Coal::CPUKernelWorkGroup::Context::context, Coal::CPUKernel::device(), Coal::Program::deviceDependentProgram(), getContextAddr(), Coal::Kernel::hasLocals(), incVec(), Coal::CPUProgram::jit(), Coal::CPUKernel::kernel(), Coal::CPUKernelWorkGroup::Context::local_id, p_args, p_current_context, p_current_work_item, p_dummy_context, p_had_barrier, p_kernel, p_kernel_func_addr, p_max_local_id, p_num_work_items, p_work_dim, Coal::Object::parent(), and setThreadLocalWorkGroup().
Referenced by worker().
void CPUKernelWorkGroup::writeImage | ( | Image2D * | image, |
int | x, | ||
int | y, | ||
int | z, | ||
int32_t * | color | ||
) | const |
Definition at line 460 of file cpu/sampler.cpp.
void CPUKernelWorkGroup::writeImage | ( | Image2D * | image, |
int | x, | ||
int | y, | ||
int | z, | ||
uint32_t * | color | ||
) | const |
Definition at line 466 of file cpu/sampler.cpp.
void CPUKernelWorkGroup::writeImage | ( | Image2D * | image, |
int | x, | ||
int | y, | ||
int | z, | ||
float * | color | ||
) | const |
Definition at line 454 of file cpu/sampler.cpp.
Referenced by write_imagef(), write_imagei(), and write_imageui().
void CPUKernelWorkGroup::writeImageImpl | ( | Image2D * | image, |
int | x, | ||
int | y, | ||
int | z, | ||
T * | color | ||
) | const [private] |
Definition at line 434 of file cpu/sampler.cpp.
References Coal::Image2D::channels(), convert_to_format(), Coal::Image2D::format(), getImageData(), and swizzle().
void* Coal::CPUKernelWorkGroup::p_args [private] |
Definition at line 265 of file cpu/kernel.h.
void* Coal::CPUKernelWorkGroup::p_contexts [private] |
Definition at line 279 of file cpu/kernel.h.
Referenced by barrier(), and getContextAddr().
Definition at line 257 of file cpu/kernel.h.
Referenced by callArgs(), and ~CPUKernelWorkGroup().
Definition at line 277 of file cpu/kernel.h.
Referenced by barrier(), getGlobalId(), getLocalID(), and run().
unsigned int Coal::CPUKernelWorkGroup::p_current_work_item [private] |
Definition at line 281 of file cpu/kernel.h.
Definition at line 278 of file cpu/kernel.h.
Referenced by run().
KernelEvent* Coal::CPUKernelWorkGroup::p_event [private] |
Definition at line 258 of file cpu/kernel.h.
Referenced by getGlobalOffset(), getGlobalSize(), getLocalSize(), and getNumGroups().
size_t Coal::CPUKernelWorkGroup::p_global_id_start_offset[MAX_WORK_DIMS] [private] |
Definition at line 260 of file cpu/kernel.h.
Referenced by CPUKernelWorkGroup(), and getGlobalId().
bool Coal::CPUKernelWorkGroup::p_had_barrier [private] |
Definition at line 282 of file cpu/kernel.h.
size_t Coal::CPUKernelWorkGroup::p_index[MAX_WORK_DIMS] [private] |
Definition at line 260 of file cpu/kernel.h.
Referenced by CPUKernelWorkGroup(), and getGroupID().
CPUKernel* Coal::CPUKernelWorkGroup::p_kernel [private] |
Definition at line 256 of file cpu/kernel.h.
Referenced by barrier(), builtinNotFound(), callArgs(), getImageData(), and run().
void(* Coal::CPUKernelWorkGroup::p_kernel_func_addr)(void *) [private] |
Definition at line 264 of file cpu/kernel.h.
size_t Coal::CPUKernelWorkGroup::p_max_local_id[MAX_WORK_DIMS] [private] |
Definition at line 260 of file cpu/kernel.h.
Referenced by barrier(), CPUKernelWorkGroup(), and run().
unsigned int Coal::CPUKernelWorkGroup::p_num_work_items [private] |
Definition at line 281 of file cpu/kernel.h.
Referenced by barrier(), CPUKernelWorkGroup(), and run().
size_t Coal::CPUKernelWorkGroup::p_stack_size [private] |
Definition at line 280 of file cpu/kernel.h.
Referenced by barrier(), and getContextAddr().
cl_uint Coal::CPUKernelWorkGroup::p_work_dim [private] |
Definition at line 259 of file cpu/kernel.h.
Referenced by barrier(), CPUKernelWorkGroup(), getGlobalId(), getGlobalOffset(), getGlobalSize(), getGroupID(), getLocalID(), getLocalSize(), getNumGroups(), getWorkDim(), and run().