TOpenCLKernel.Enqueue Method

Overload List

#SignatureDescription
1procedure Enqueue(const CmdQueue: TOpenCLCommandQueue; WorkSize: Integer; CPUAdjust: Boolean);Submits the Kernel to cmdQueue for computation with specified WorkSize.
2procedure Enqueue(const CmdQueue: TOpenCLCommandQueue; WorkSize: Integer; LocalSize: Integer; CPUAdjust: Boolean);Submits the Kernel to cmdQueue for computation with specified WorkSize and LocalSize.

Overload 1: procedure Enqueue(const CmdQueue: TOpenCLCommandQueue; WorkSize: Integer; CPUAdjust: Boolean);

Submits the Kernel to cmdQueue for computation with specified WorkSize.

#NameTypeDescription
1CmdQueueTOpenCLCommandQueue
2WorkSizeInteger
3CPUAdjustBoolean

Result: stored in self (calling object)

Remarks:

Setting CPUAdjust to true will reduce WorkSize by factor OPENCL_BLOCKLEN and assume presence of kernel internal for-loops. Kernel internal for-loops can significantly speed up execution of the kernel on CPU devices lowering the function call overhead. The CPUAdjust parameter is used only if the device is of CPU type. The for-loop pattern expected inside the kernel looks like this:

  size_t i = get_global_id(0);
	size_t tIdx = i*BLOCK_LEN;
	size_t tIdxLen = tIdx + BLOCK_LEN;
	if (tIdxLen > Length) tIdxLen = Length;
	for (i = tIdx; i < tIdxLen; i++)
	{

  }

where BLOCK_LEN matches OPENCL_BLOCKLEN.

Overload 2: procedure Enqueue(const CmdQueue: TOpenCLCommandQueue; WorkSize: Integer; LocalSize: Integer; CPUAdjust: Boolean);

Submits the Kernel to cmdQueue for computation with specified WorkSize and LocalSize.

#NameTypeDescription
1CmdQueueTOpenCLCommandQueue
2WorkSizeInteger
3LocalSizeInteger
4CPUAdjustBoolean

Result: stored in self (calling object)

Remarks:

When specified explicitely the (WorkSize mod LocalSize) is required to be zero. Setting CPUAdjust to true will reduce WorkSize by factor OPENCL_BLOCKLEN and assume presence of kernel internal for-loops. Kernel internal for-loops can significantly speed up execution of the kernel on CPU devices lowering the function call overhead. Kernel internal for-loops in GPU devices cause large performance penalties. The CPUAdjust parameter should be used only if the device is of CPU type. LocalSize is also called workgroup size. The for-loop pattern expected inside the kernel looks like this:

  size_t i = get_global_id(0);
	size_t tIdx = i*BLOCK_LEN;
	size_t tIdxLen = tIdx + BLOCK_LEN;
	if (tIdxLen > Length) tIdxLen = Length;
	for (i = tIdx; i < tIdxLen; i++)
	{

  }

where BLOCK_LEN matches OPENCL_BLOCKLEN.