How AFOCL's pocl almaif driver slices tasks and schedules them when processing NDRange ?

[soccercheng]

Hi AFOCL Team

I've been trying to add my first OpenCL kernel (written in OpenCL C), and I have the following questions.

• Does AFOCL's slice tasks and schedules them when processing NDRange ?
• How does AFOCL handle OpenCL build-in functions, like "size_t get_group_id(uint dimidx)", "size_t get_global_id(uint dimidx)" or size_t get_local_id(uint dimidx)
• How AFOCL's pocl almaif driver schedules NDRange ? It seems that current fields of "struct AQLDispatchPacket" does NOT have enough information to re-construct the necessary input to the HLS'ed OpenCL kernel IP.

~Sting

[TopiLeppanen]

1. No. The built-in kernel implementation is responsible for executing the entire NDRange
2. It doesn't. Those are OpenCL C kernel language functions. AFOCL specifically works with built-in kernels, so the kernel implementation is free to do whatever to execute the NDRange.
3. The entire NDRange is given in the single AQLDispatchPacket. Then the firmware+HLS'ed component can together execute that however they want. What do you mean "re-construct the necessary input"? Do you mean the input arguments, or something else? The current implementation doesn't really use work-group abstraction at all IIRC, since it doesn't make much of a difference when executing e.g. vector addition. The current implementation always executes the entire NDRange however it wants. E.g. sequentially (with some vectorization factor).

In the future, it might make sense to scale up the FPGA utilization by having multiple HLS'ed components exist in parallel which are fed work-groups to execute by the firmware. Scaling up by simple vectorization gets quite difficult with hundreds of elements (FPGA gets congested, or is not able to meet the timing)

[soccercheng]

As for the 3rd question, what we're trying to do is to extend current AFOCL built-in kernels by adding a new ccccccccccwritten in "OpenCL C", which uses OpenCL build-in functions, like get_global_id... since we didn't know that there is no NDRange scheduler in current AFOCL firmware implementation.

What I meant was, from current "struct AQLDispatchPacket", we're not able to re-construct the necessary input to control registers of new HLS'ed kernel written in OpenCL C.

Current AQLDispatchPacket contains only dimensions, workgroup_size_x and grid_size_x, and the control registers of newly HSL'ed kernel require the following fields:

local_size_x :out STD_LOGIC_VECTOR(31 downto 0);
local_size_y :out STD_LOGIC_VECTOR(31 downto 0);
local_size_z :out STD_LOGIC_VECTOR(31 downto 0);
group_id_x :out STD_LOGIC_VECTOR(31 downto 0);
group_id_y :out STD_LOGIC_VECTOR(31 downto 0);
group_id_z :out STD_LOGIC_VECTOR(31 downto 0);
global_offset_x :out STD_LOGIC_VECTOR(31 downto 0);
global_offset_y :out STD_LOGIC_VECTOR(31 downto 0);
global_offset_z :out STD_LOGIC_VECTOR(31 downto 0);

[TopiLeppanen]

Right, I see. There is a pocl_context(_32b) struct that you can take a look at. You might need to make some changes to get that information all the way through