On 02/05/2014 03:57 PM, Panagiotis Apostolou wrote:
> As far as i have understood call_pocl_workgroup() cannot be skipped as
> it links the kernel with the kernel built-in functions
Yes, the hook should be added to the call to LLVM.
I.e. convert this
POCL_LOCK(kernel_compiler_lock);
// TODO pass these as parameters instead, this is not thread safe!
pocl::LocalSize.clear();
pocl::LocalSize.addValue(local_x);
pocl::LocalSize.addValue(local_y);
pocl::LocalSize.addValue(local_z);
KernelName = kernel->name;
kernel_compiler_passes(device, linked_bc->getDataLayout()).run(*linked_bc);
POCL_UNLOCK(kernel_compiler_lock);
to something like:
POCL_LOCK(kernel_compiler_lock);
if (device->ops->workgroup_compiler !=)
device->ops->work_group_compiler(
device->data, linked_bc, local_x, local_y, local_z);
else
default_work_group_compiler(
device->data, linked_bc, local_x, local_y, local_z);
POCL_UNLOCK(kernel_compiler_lock);
That is, in case the device does not want to override the work-group
generation part, it falls back to the current behavior that generates
a function with all the work-items executed with one call.
> So when
> call_pocl_workgroup(command_queue->device, kernel, local_x, local_y,
> local_z, ...)
> is called form the clEnqueueNDRangeKernel(), should i replace the
> local_x, y and z values with something
> that my driver finds appropriate (for example 1)?
Those are the dimensions of the work-group. You might not need them
in your override as you do not generate a work-group function, but
they should appear in the hook API so one can use them if needed.
--
Pekka
------------------------------------------------------------------------------
Managing the Performance of Cloud-Based Applications
Take advantage of what the Cloud has to offer - Avoid Common Pitfalls.
Read the Whitepaper.
http://pubads.g.doubleclick.net/gampad/clk?id=121051231&iu=/4140/ostg.clktrk
_______________________________________________
pocl-devel mailing list
[email protected]
https://lists.sourceforge.net/lists/listinfo/pocl-devel
