Device Global Copy Kernel Support

[sherry-yuan]

Theoretically it should be (given the lazy programming feature is available in runtime), but need to double check / at least layout how things should be called.

The precise questions are:

1. Are there any issues in programming the FPGA, given a cl_program object?
2. Any other foreseeable issues with launching the kernel from within the Runtime?
3. What other information should the compiler provide through the autodiscovery string (now there is mapping of device_global name to pointer, all properties associated with device global)

CC: @zibaiwan @pcolberg @aditikum

[sherry-yuan]

Short Answers

1. Are there any issues in programming the FPGA, given a cl_program object? (No issue)
2. Any other foreseeable issues with launching the kernel from within the Runtime? (No issue that has no workaround as of now)
3. What other information should the compiler provide through the autodiscovery string (now there is mapping of device_global name to pointer, all properties associated with device global) (see considerations below for things needed when launching a kernel).

Currently Drafted Solution

1. Probably need to program, which we can borrow the logic from clCreateProgramWithBinaryAndProgramDeviceIntelFPGA (there should be no issue)

   fpga-runtime-for-opencl/src/acl_program.cpp

   Lines 544 to 597 in 950f21d

   	for (idev = 0; idev < program->num_devices; idev++) {
   	acl_device_program_info_t *dev_prog = program->dev_prog[idev];
   	if (dev_prog && dev_prog->build_status == CL_BUILD_SUCCESS) {
   	cl_int result;
   	// Use a regular event with an internal command type.
   	// We want a regular event so we get profiling info, so that we
   	// have the infrastructure to measure delays, and the user
   	// could in principle optimize this delay away.
   	cl_event reprogram_event = 0;
   	// Just use the auto_queue. We really only support one device
   	// anyway.
   	cl_command_queue cq = program->context->auto_queue;
   	// Schedule an eager programming of the device.
   	acl_print_debug_msg(
   	"Device is not yet programmed: plan to eagerly program it\n");
   	result = acl_create_event(cq, 0, 0, // Don't wait on other events
   	CL_COMMAND_PROGRAM_DEVICE_INTELFPGA,
   	&reprogram_event);
   	if (result == CL_SUCCESS) {
   	acl_device_op_t reprogram_op;
   	reprogram_event->cmd.info.eager_program = &(dev_prog->device_binary);
   	// Try scheduling it.
   	reprogram_op.link = ACL_OPEN;
   	acl_device_op_reset_device_op(&reprogram_op);
   	reprogram_op.status = ACL_PROPOSED;
   	reprogram_op.execution_status = ACL_PROPOSED;
   	reprogram_op.info.type = ACL_DEVICE_OP_REPROGRAM;
   	reprogram_op.info.event = reprogram_event;
   	reprogram_op.info.index = 0;
   	reprogram_op.conflict_type = acl_device_op_conflict_type(&reprogram_op);
   	acl_program_device(NULL, &reprogram_op);
   	if (reprogram_op.execution_status != CL_SUCCESS) {
   	UNLOCK_BAIL_INFO(CL_DEVICE_NOT_AVAILABLE, context,
   	"Reprogram of device failed");
   	}
   	} else {
   	UNLOCK_BAIL_INFO(CL_OUT_OF_HOST_MEMORY, context, "Invalid binary");
   	}
   	} else {
   	UNLOCK_BAIL_INFO(CL_BUILD_PROGRAM_FAILURE, context,
   	"Program is not built correctly");
   	}
   	}

2. clCreateKernelIntelFPGA(program, name, errcode_ret)
3. extract dest pointer from autodiscovery
4. clSetKernelArgIntelFPGA (need to set src, dest, offset, size)
5. clEnqueueTask
6. Depending on whether blocking: wait/ not wait for event_wait_list

Considerations

1. [This is impl detail, feel free to skip] Make sure there is no double lock when calling other opencl function, otherwise it will hung.
2. Do we know if clEnqueueWirteBuffer is called before passing buffer into the function. (although I am not too worried about it given it will be enqueued anyways when enqueue kernel: l_copy_and_adjust_arguments_for_device)
3. The runtime is not going to know the device memory interface through the platform global var without receiving aocx. That's at least the case for simulation flow.
4. Above potentially bring a problem when we need to allocate buffer for the buffer that is passed in (and especially when that buffer is targeted at multi-memory system, since memory operation do not have direct access to autodiscovery), depending on whether buffer allocation is deferred.
5. What was the consideration of not reusing l_enqueue_mem_transfer?

   fpga-runtime-for-opencl/src/acl_mem.cpp

   Lines 4706 to 5154 in d9df7a9

   	cl_int l_enqueue_mem_transfer(cl_command_queue command_queue, cl_bool blocking,
   	cl_mem src_buffer, size_t src_offset[3],
   	size_t src_row_pitch, size_t src_slice_pitch,
   	cl_mem dst_buffer, size_t dst_offset[3],
   	size_t dst_row_pitch, size_t dst_slice_pitch,
   	size_t cb[3], cl_uint num_events,
   	const cl_event *events, cl_event *event,
   	cl_command_type type, cl_map_flags map_flags) {
   	acl_assert_locked();
   	if (!acl_command_queue_is_valid(command_queue)) {
   	return CL_INVALID_COMMAND_QUEUE;
   	}
   	cl_context context = command_queue->context;
   	if ((num_events > 0 && !events) || (!num_events && events)) {
   	ERR_RET(CL_INVALID_EVENT_WAIT_LIST, context,
   	"Invalid event_wait_list or num_events_in_wait_list");
   	}
   	if (!acl_mem_is_valid(src_buffer))
   	ERR_RET(CL_INVALID_MEM_OBJECT, context, "Source buffer is invalid");
   	if (!acl_mem_is_valid(dst_buffer))
   	ERR_RET(CL_INVALID_MEM_OBJECT, context, "Destination buffer is invalid");
   	if (command_queue->context != src_buffer->context)
   	ERR_RET(CL_INVALID_CONTEXT, context,
   	"Source buffer is not in the same context as the command queue");
   	if (command_queue->context != dst_buffer->context)
   	ERR_RET(
   	CL_INVALID_CONTEXT, context,
   	"Destination buffer is not in the same context as the command queue");
   	if (cb[0] == 0)
   	ERR_RET(CL_INVALID_VALUE, context, "Region 'x' size is zero");
   	if (cb[1] == 0)
   	ERR_RET(CL_INVALID_VALUE, context, "Region 'y' size is zero");
   	if (cb[2] == 0)
   	ERR_RET(CL_INVALID_VALUE, context, "Region 'z' size is zero");
   	// Check that the requested operation is compatible with the flags the src/dst
   	// buffers were created with. Checks for image related operations have not
   	// been implemented.
   	switch (type) {
   	case CL_COMMAND_READ_BUFFER:
   	case CL_COMMAND_READ_BUFFER_RECT:
   	if (src_buffer->flags & CL_MEM_HOST_WRITE_ONLY ||
   	src_buffer->flags & CL_MEM_HOST_NO_ACCESS) {
   	ERR_RET(CL_INVALID_OPERATION, context,
   	"clEnqeueueReadBuffer cannot be called on a buffer "
   	"created with CL_MEM_HOST_WRITE_ONLY or CL_MEM_HOST_NO_ACCESS");
   	}
   	break;
   	case CL_COMMAND_WRITE_BUFFER:
   	case CL_COMMAND_WRITE_BUFFER_RECT:
   	if (dst_buffer->flags & CL_MEM_COPY_HOST_PTR &&
   	dst_buffer->copy_host_ptr_skip_check) {
   	// Don't block the call to clEnqueueWriteBuffer when we're initializing
   	// a read-only or no-access buffer created with CL_MEM_COPY_HOST_PTR since
   	// it's an internal usage of the API.
   	dst_buffer->copy_host_ptr_skip_check = CL_FALSE;
   	break;
   	}
   	if (dst_buffer->flags & CL_MEM_HOST_READ_ONLY ||
   	dst_buffer->flags & CL_MEM_HOST_NO_ACCESS) {
   	ERR_RET(CL_INVALID_OPERATION, context,
   	"clEnqeueueWriteBuffer cannot be called on a buffer "
   	"created with CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_NO_ACCESS");
   	}
   	break;
   	case CL_COMMAND_COPY_BUFFER:
   	case CL_COMMAND_COPY_BUFFER_RECT:
   	case CL_COMMAND_FILL_BUFFER:
   	case CL_COMMAND_UNMAP_MEM_OBJECT:
   	break;
   	case CL_COMMAND_MAP_BUFFER:
   	if ((src_buffer->flags & CL_MEM_HOST_WRITE_ONLY ||
   	src_buffer->flags & CL_MEM_HOST_NO_ACCESS) &&
   	map_flags & CL_MAP_READ) {
   	ERR_RET(
   	CL_INVALID_OPERATION, context,
   	"clEnqeueueMapBuffer with CL_MAP_READ cannot be called on a buffer "
   	"created with CL_MEM_HOST_WRITE_ONLY or CL_MEM_HOST_NO_ACCESS");
   	}
   	if ((src_buffer->flags & CL_MEM_HOST_READ_ONLY ||
   	src_buffer->flags & CL_MEM_HOST_NO_ACCESS) &&
   	(map_flags & CL_MAP_WRITE ||
   	map_flags & CL_MAP_WRITE_INVALIDATE_REGION)) {
   	ERR_RET(CL_INVALID_OPERATION, context,
   	"clEnqeueueMapBuffer with CL_MAP_WRITE or "
   	"CL_MAP_WRITE_INVALIDATE_REGION "
   	"cannot be called on a buffer created with CL_MEM_HOST_READ_ONLY "
   	"or CL_MEM_HOST_NO_ACCESS");
   	}
   	break;
   	case CL_COMMAND_READ_IMAGE:
   	case CL_COMMAND_WRITE_IMAGE:
   	case CL_COMMAND_COPY_IMAGE:
   	case CL_COMMAND_COPY_IMAGE_TO_BUFFER:
   	case CL_COMMAND_COPY_BUFFER_TO_IMAGE:
   	case CL_COMMAND_MAP_IMAGE:
   	case CL_COMMAND_FILL_IMAGE:
   	#ifndef ACL_SUPPORT_IMAGES
   	ERR_RET(CL_INVALID_OPERATION, context, "Device does not support images");
   	#endif
   	break;
   	default:
   	assert(0 && "Command is not a memory transfer related command type");
   	}
   	cl_int errcode_ret;
   	size_t src_element_size;
   	// Check that the copy area of the source buffer is within the allocated
   	// memory
   	switch (src_buffer->mem_object_type) {
   	size_t default_slice_pitch;
   	case CL_MEM_OBJECT_BUFFER:
   	if ((src_offset[2] + cb[2] - 1) * src_slice_pitch +
   	(src_offset[1] + cb[1] - 1) * src_row_pitch + src_offset[0] +
   	cb[0] >
   	src_buffer->size)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Source buffer offset plus byte count exceeds source buffer size");
   	if (src_row_pitch != 0 && src_row_pitch < cb[0])
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source buffer row pitch is less than region 'x' size");
   	if (dst_buffer->mem_object_type == CL_MEM_OBJECT_IMAGE1D_ARRAY) {
   	default_slice_pitch = src_row_pitch;
   	} else {
   	default_slice_pitch = cb[1] * src_row_pitch;
   	}
   	if (src_slice_pitch != 0 && src_slice_pitch < default_slice_pitch) {
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Source buffer slice pitch is less than region 'y' size * row pitch");
   	}
   	break;
   	case CL_MEM_OBJECT_IMAGE1D:
   	case CL_MEM_OBJECT_IMAGE1D_BUFFER:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] != 0 || src_offset[2] != 0)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in source image");
   	if (cb[1] != 1 || cb[2] != 1)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in source image");
   	break;
   	case CL_MEM_OBJECT_IMAGE1D_ARRAY:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image array size");
   	if (src_offset[2] != 0)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in source image");
   	if (cb[2] != 1)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in source image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image height");
   	if (src_offset[2] != 0)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in source image");
   	if (cb[2] != 1)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in source image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D_ARRAY:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image height");
   	if (src_offset[2] + cb[2] >
   	src_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image array size");
   	break;
   	case CL_MEM_OBJECT_IMAGE3D:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image height");
   	if (src_offset[2] + cb[2] >
   	src_buffer->fields.image_objs.image_desc->image_depth)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image depth");
   	break;
   	default:
   	std::stringstream ss;
   	ss << "Memory transfers of source mem object type not supported 0x"
   	<< std::hex << src_buffer->mem_object_type;
   	ERR_RET(CL_INVALID_VALUE, context, ss.str().c_str());
   	break;
   	}
   	size_t dst_element_size;
   	// Check that the copy area of the destination buffer is within the allocated
   	// memory
   	switch (dst_buffer->mem_object_type) {
   	case CL_MEM_OBJECT_BUFFER:
   	if ((dst_offset[2] + cb[2] - 1) * dst_slice_pitch +
   	(dst_offset[1] + cb[1] - 1) * dst_row_pitch + dst_offset[0] +
   	cb[0] >
   	dst_buffer->size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination buffer offset plus byte count exceeds destination "
   	"buffer size");
   	if (dst_row_pitch != 0 && dst_row_pitch < cb[0])
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination buffer row pitch is less than region 'x' size");
   	if (dst_slice_pitch != 0 && dst_slice_pitch < cb[1] * dst_row_pitch)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination buffer slice pitch is less than region 'y' size * "
   	"row pitch");
   	break;
   	case CL_MEM_OBJECT_IMAGE1D:
   	case CL_MEM_OBJECT_IMAGE1D_BUFFER:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] != 0 || dst_offset[2] != 0)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in destination image");
   	if (cb[1] != 1 || cb[2] != 1)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in destination image");
   	break;
   	case CL_MEM_OBJECT_IMAGE1D_ARRAY:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image array size");
   	if (dst_offset[2] != 0)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in destination image");
   	if (cb[2] != 1)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in destination image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image height");
   	if (dst_offset[2] != 0)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in destination image");
   	if (cb[2] != 1)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in destination image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D_ARRAY:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image height");
   	if (dst_offset[2] + cb[2] >
   	dst_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image array size");
   	break;
   	case CL_MEM_OBJECT_IMAGE3D:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size) {
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	}
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_height) {
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image height");
   	}
   	if (dst_offset[2] + cb[2] >
   	dst_buffer->fields.image_objs.image_desc->image_depth) {
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image depth");
   	}
   	break;
   	default:
   	std::stringstream ss;
   	ss << "Memory transfers of destination mem object type not supported 0x"
   	<< std::hex << dst_buffer->mem_object_type;
   	ERR_RET(CL_INVALID_VALUE, context, ss.str().c_str());
   	break;
   	}
   	cl_event local_event = 0; // used for blocking
   	// Create an event/command to actually move the data at the appropriate
   	// time.
   	cl_int status =
   	acl_create_event(command_queue, num_events, events, type, &local_event);
   	if (status != CL_SUCCESS)
   	return status; // already signalled callback
   	local_event->cmd.info.mem_xfer.src_mem = src_buffer;
   	clRetainMemObject(src_buffer);
   	local_event->cmd.info.mem_xfer.dst_mem = dst_buffer;
   	clRetainMemObject(dst_buffer);
   	for (size_t i = 0; i < 3; ++i) {
   	local_event->cmd.info.mem_xfer.src_offset[i] = src_offset[i];
   	local_event->cmd.info.mem_xfer.dst_offset[i] = dst_offset[i];
   	local_event->cmd.info.mem_xfer.cb[i] = cb[i];
   	}
   	local_event->cmd.info.mem_xfer.src_row_pitch = src_row_pitch;
   	local_event->cmd.info.mem_xfer.src_slice_pitch = src_slice_pitch;
   	local_event->cmd.info.mem_xfer.dst_row_pitch = dst_row_pitch;
   	local_event->cmd.info.mem_xfer.dst_slice_pitch = dst_slice_pitch;
   	local_event->cmd.info.mem_xfer.map_flags = map_flags;
   	local_event->cmd.info.mem_xfer.is_auto_map = 0;
   	acl_idle_update(
   	command_queue
   	->context); // If nothing's blocking, then complete right away
   	if (blocking) {
   	status = clWaitForEvents(1, &local_event);
   	}
   	if (event) {
   	*event = local_event;
   	} else {
   	// User didn't care, so forget about the event.
   	clReleaseEvent(local_event);
   	acl_idle_update(command_queue->context); // Clean up early
   	}
   	if (blocking && status == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST)
   	return status;
   	return CL_SUCCESS;
   	}

6. If buffer alloc is deferred / not yet allocated, then enqueueing this kernel will result in an extra copy that we can avoid. i.e we can directly copy from buffer's host to device_global rather than host -> device -> device global. But maybe we don't care about performance.

Resources

1. Implementation Doc
2. OpenCL Spec
3. Sycl Spec

[sherry-yuan]

Summary

This is the summary version of above comment, see above for more detail. Feel free to comment if there is anything missed.

Short Answers

1. Are there any issues in programming the FPGA, given a cl_program object? (No issue)
2. Any other foreseeable issues with launching the kernel from within the Runtime? (No issue as of now)
3. What other information should the compiler provide through the autodiscovery string (now there is mapping of device_global name to pointer, all properties associated with device global) (this probably doesn't matter for mem copy kernel: launching a kernel also requires arguments for workgroup size etc).

Launch kernel arguments: l_enqueue_kernel_with_type(command_queue, kernel, ?work_dim, ?global_work_offset, ?global_work_size, ?local_work_size, num_events_in_wait_list, event_wait_list, event, ?CL_COMMAND_MIGRATE_MEM_OBJECTS);

Always pass in constant for size and offsets?

Main Questions

1. What was the consideration of not reusing l_enqueue_mem_transfer (is it because dest is not buffer)?

   fpga-runtime-for-opencl/src/acl_mem.cpp

   Lines 4706 to 5154 in d9df7a9

   	cl_int l_enqueue_mem_transfer(cl_command_queue command_queue, cl_bool blocking,
   	cl_mem src_buffer, size_t src_offset[3],
   	size_t src_row_pitch, size_t src_slice_pitch,
   	cl_mem dst_buffer, size_t dst_offset[3],
   	size_t dst_row_pitch, size_t dst_slice_pitch,
   	size_t cb[3], cl_uint num_events,
   	const cl_event *events, cl_event *event,
   	cl_command_type type, cl_map_flags map_flags) {
   	acl_assert_locked();
   	if (!acl_command_queue_is_valid(command_queue)) {
   	return CL_INVALID_COMMAND_QUEUE;
   	}
   	cl_context context = command_queue->context;
   	if ((num_events > 0 && !events) || (!num_events && events)) {
   	ERR_RET(CL_INVALID_EVENT_WAIT_LIST, context,
   	"Invalid event_wait_list or num_events_in_wait_list");
   	}
   	if (!acl_mem_is_valid(src_buffer))
   	ERR_RET(CL_INVALID_MEM_OBJECT, context, "Source buffer is invalid");
   	if (!acl_mem_is_valid(dst_buffer))
   	ERR_RET(CL_INVALID_MEM_OBJECT, context, "Destination buffer is invalid");
   	if (command_queue->context != src_buffer->context)
   	ERR_RET(CL_INVALID_CONTEXT, context,
   	"Source buffer is not in the same context as the command queue");
   	if (command_queue->context != dst_buffer->context)
   	ERR_RET(
   	CL_INVALID_CONTEXT, context,
   	"Destination buffer is not in the same context as the command queue");
   	if (cb[0] == 0)
   	ERR_RET(CL_INVALID_VALUE, context, "Region 'x' size is zero");
   	if (cb[1] == 0)
   	ERR_RET(CL_INVALID_VALUE, context, "Region 'y' size is zero");
   	if (cb[2] == 0)
   	ERR_RET(CL_INVALID_VALUE, context, "Region 'z' size is zero");
   	// Check that the requested operation is compatible with the flags the src/dst
   	// buffers were created with. Checks for image related operations have not
   	// been implemented.
   	switch (type) {
   	case CL_COMMAND_READ_BUFFER:
   	case CL_COMMAND_READ_BUFFER_RECT:
   	if (src_buffer->flags & CL_MEM_HOST_WRITE_ONLY ||
   	src_buffer->flags & CL_MEM_HOST_NO_ACCESS) {
   	ERR_RET(CL_INVALID_OPERATION, context,
   	"clEnqeueueReadBuffer cannot be called on a buffer "
   	"created with CL_MEM_HOST_WRITE_ONLY or CL_MEM_HOST_NO_ACCESS");
   	}
   	break;
   	case CL_COMMAND_WRITE_BUFFER:
   	case CL_COMMAND_WRITE_BUFFER_RECT:
   	if (dst_buffer->flags & CL_MEM_COPY_HOST_PTR &&
   	dst_buffer->copy_host_ptr_skip_check) {
   	// Don't block the call to clEnqueueWriteBuffer when we're initializing
   	// a read-only or no-access buffer created with CL_MEM_COPY_HOST_PTR since
   	// it's an internal usage of the API.
   	dst_buffer->copy_host_ptr_skip_check = CL_FALSE;
   	break;
   	}
   	if (dst_buffer->flags & CL_MEM_HOST_READ_ONLY ||
   	dst_buffer->flags & CL_MEM_HOST_NO_ACCESS) {
   	ERR_RET(CL_INVALID_OPERATION, context,
   	"clEnqeueueWriteBuffer cannot be called on a buffer "
   	"created with CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_NO_ACCESS");
   	}
   	break;
   	case CL_COMMAND_COPY_BUFFER:
   	case CL_COMMAND_COPY_BUFFER_RECT:
   	case CL_COMMAND_FILL_BUFFER:
   	case CL_COMMAND_UNMAP_MEM_OBJECT:
   	break;
   	case CL_COMMAND_MAP_BUFFER:
   	if ((src_buffer->flags & CL_MEM_HOST_WRITE_ONLY ||
   	src_buffer->flags & CL_MEM_HOST_NO_ACCESS) &&
   	map_flags & CL_MAP_READ) {
   	ERR_RET(
   	CL_INVALID_OPERATION, context,
   	"clEnqeueueMapBuffer with CL_MAP_READ cannot be called on a buffer "
   	"created with CL_MEM_HOST_WRITE_ONLY or CL_MEM_HOST_NO_ACCESS");
   	}
   	if ((src_buffer->flags & CL_MEM_HOST_READ_ONLY ||
   	src_buffer->flags & CL_MEM_HOST_NO_ACCESS) &&
   	(map_flags & CL_MAP_WRITE ||
   	map_flags & CL_MAP_WRITE_INVALIDATE_REGION)) {
   	ERR_RET(CL_INVALID_OPERATION, context,
   	"clEnqeueueMapBuffer with CL_MAP_WRITE or "
   	"CL_MAP_WRITE_INVALIDATE_REGION "
   	"cannot be called on a buffer created with CL_MEM_HOST_READ_ONLY "
   	"or CL_MEM_HOST_NO_ACCESS");
   	}
   	break;
   	case CL_COMMAND_READ_IMAGE:
   	case CL_COMMAND_WRITE_IMAGE:
   	case CL_COMMAND_COPY_IMAGE:
   	case CL_COMMAND_COPY_IMAGE_TO_BUFFER:
   	case CL_COMMAND_COPY_BUFFER_TO_IMAGE:
   	case CL_COMMAND_MAP_IMAGE:
   	case CL_COMMAND_FILL_IMAGE:
   	#ifndef ACL_SUPPORT_IMAGES
   	ERR_RET(CL_INVALID_OPERATION, context, "Device does not support images");
   	#endif
   	break;
   	default:
   	assert(0 && "Command is not a memory transfer related command type");
   	}
   	cl_int errcode_ret;
   	size_t src_element_size;
   	// Check that the copy area of the source buffer is within the allocated
   	// memory
   	switch (src_buffer->mem_object_type) {
   	size_t default_slice_pitch;
   	case CL_MEM_OBJECT_BUFFER:
   	if ((src_offset[2] + cb[2] - 1) * src_slice_pitch +
   	(src_offset[1] + cb[1] - 1) * src_row_pitch + src_offset[0] +
   	cb[0] >
   	src_buffer->size)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Source buffer offset plus byte count exceeds source buffer size");
   	if (src_row_pitch != 0 && src_row_pitch < cb[0])
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source buffer row pitch is less than region 'x' size");
   	if (dst_buffer->mem_object_type == CL_MEM_OBJECT_IMAGE1D_ARRAY) {
   	default_slice_pitch = src_row_pitch;
   	} else {
   	default_slice_pitch = cb[1] * src_row_pitch;
   	}
   	if (src_slice_pitch != 0 && src_slice_pitch < default_slice_pitch) {
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Source buffer slice pitch is less than region 'y' size * row pitch");
   	}
   	break;
   	case CL_MEM_OBJECT_IMAGE1D:
   	case CL_MEM_OBJECT_IMAGE1D_BUFFER:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] != 0 || src_offset[2] != 0)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in source image");
   	if (cb[1] != 1 || cb[2] != 1)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in source image");
   	break;
   	case CL_MEM_OBJECT_IMAGE1D_ARRAY:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image array size");
   	if (src_offset[2] != 0)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in source image");
   	if (cb[2] != 1)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in source image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image height");
   	if (src_offset[2] != 0)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in source image");
   	if (cb[2] != 1)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in source image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D_ARRAY:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image height");
   	if (src_offset[2] + cb[2] >
   	src_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image array size");
   	break;
   	case CL_MEM_OBJECT_IMAGE3D:
   	src_element_size = acl_get_image_element_size(
   	src_buffer->context, src_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (src_offset[0] + cb[0] >
   	src_buffer->fields.image_objs.image_desc->image_width *
   	src_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image width");
   	if (src_offset[1] + cb[1] >
   	src_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image height");
   	if (src_offset[2] + cb[2] >
   	src_buffer->fields.image_objs.image_desc->image_depth)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Source image offset plus region exceeds image depth");
   	break;
   	default:
   	std::stringstream ss;
   	ss << "Memory transfers of source mem object type not supported 0x"
   	<< std::hex << src_buffer->mem_object_type;
   	ERR_RET(CL_INVALID_VALUE, context, ss.str().c_str());
   	break;
   	}
   	size_t dst_element_size;
   	// Check that the copy area of the destination buffer is within the allocated
   	// memory
   	switch (dst_buffer->mem_object_type) {
   	case CL_MEM_OBJECT_BUFFER:
   	if ((dst_offset[2] + cb[2] - 1) * dst_slice_pitch +
   	(dst_offset[1] + cb[1] - 1) * dst_row_pitch + dst_offset[0] +
   	cb[0] >
   	dst_buffer->size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination buffer offset plus byte count exceeds destination "
   	"buffer size");
   	if (dst_row_pitch != 0 && dst_row_pitch < cb[0])
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination buffer row pitch is less than region 'x' size");
   	if (dst_slice_pitch != 0 && dst_slice_pitch < cb[1] * dst_row_pitch)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination buffer slice pitch is less than region 'y' size * "
   	"row pitch");
   	break;
   	case CL_MEM_OBJECT_IMAGE1D:
   	case CL_MEM_OBJECT_IMAGE1D_BUFFER:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] != 0 || dst_offset[2] != 0)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in destination image");
   	if (cb[1] != 1 || cb[2] != 1)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in destination image");
   	break;
   	case CL_MEM_OBJECT_IMAGE1D_ARRAY:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image array size");
   	if (dst_offset[2] != 0)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in destination image");
   	if (cb[2] != 1)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in destination image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image height");
   	if (dst_offset[2] != 0)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-zero offset specified for invalid index in destination image");
   	if (cb[2] != 1)
   	ERR_RET(
   	CL_INVALID_VALUE, context,
   	"Non-one region specified for invalid index in destination image");
   	break;
   	case CL_MEM_OBJECT_IMAGE2D_ARRAY:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_height)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image height");
   	if (dst_offset[2] + cb[2] >
   	dst_buffer->fields.image_objs.image_desc->image_array_size)
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image array size");
   	break;
   	case CL_MEM_OBJECT_IMAGE3D:
   	dst_element_size = acl_get_image_element_size(
   	dst_buffer->context, dst_buffer->fields.image_objs.image_format,
   	&errcode_ret);
   	if (errcode_ret != CL_SUCCESS) {
   	return errcode_ret;
   	}
   	if (dst_offset[0] + cb[0] >
   	dst_buffer->fields.image_objs.image_desc->image_width *
   	dst_element_size) {
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image width");
   	}
   	if (dst_offset[1] + cb[1] >
   	dst_buffer->fields.image_objs.image_desc->image_height) {
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image height");
   	}
   	if (dst_offset[2] + cb[2] >
   	dst_buffer->fields.image_objs.image_desc->image_depth) {
   	ERR_RET(CL_INVALID_VALUE, context,
   	"Destination image offset plus region exceeds image depth");
   	}
   	break;
   	default:
   	std::stringstream ss;
   	ss << "Memory transfers of destination mem object type not supported 0x"
   	<< std::hex << dst_buffer->mem_object_type;
   	ERR_RET(CL_INVALID_VALUE, context, ss.str().c_str());
   	break;
   	}
   	cl_event local_event = 0; // used for blocking
   	// Create an event/command to actually move the data at the appropriate
   	// time.
   	cl_int status =
   	acl_create_event(command_queue, num_events, events, type, &local_event);
   	if (status != CL_SUCCESS)
   	return status; // already signalled callback
   	local_event->cmd.info.mem_xfer.src_mem = src_buffer;
   	clRetainMemObject(src_buffer);
   	local_event->cmd.info.mem_xfer.dst_mem = dst_buffer;
   	clRetainMemObject(dst_buffer);
   	for (size_t i = 0; i < 3; ++i) {
   	local_event->cmd.info.mem_xfer.src_offset[i] = src_offset[i];
   	local_event->cmd.info.mem_xfer.dst_offset[i] = dst_offset[i];
   	local_event->cmd.info.mem_xfer.cb[i] = cb[i];
   	}
   	local_event->cmd.info.mem_xfer.src_row_pitch = src_row_pitch;
   	local_event->cmd.info.mem_xfer.src_slice_pitch = src_slice_pitch;
   	local_event->cmd.info.mem_xfer.dst_row_pitch = dst_row_pitch;
   	local_event->cmd.info.mem_xfer.dst_slice_pitch = dst_slice_pitch;
   	local_event->cmd.info.mem_xfer.map_flags = map_flags;
   	local_event->cmd.info.mem_xfer.is_auto_map = 0;
   	acl_idle_update(
   	command_queue
   	->context); // If nothing's blocking, then complete right away
   	if (blocking) {
   	status = clWaitForEvents(1, &local_event);
   	}
   	if (event) {
   	*event = local_event;
   	} else {
   	// User didn't care, so forget about the event.
   	clReleaseEvent(local_event);
   	acl_idle_update(command_queue->context); // Clean up early
   	}
   	if (blocking && status == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST)
   	return status;
   	return CL_SUCCESS;
   	}

2. If buffer alloc is deferred / not yet allocated, then enqueueing this kernel will result in an extra copy that we can avoid. i.e we can directly copy from buffer's host to device_global rather than host -> device -> device global. But maybe we don't care about performance.

[sherry-yuan]

Edit: Discard this comment as there is a better solution below provided by Artem.

In regards to the question of how work_size (both local and global should be decided).

1. If workgroup size is known at compile time instead of run time, simpler hardware is generated.
2. work group size is fixed, so number of work groups actually determines the total amount of work. This is limited by the size of each global memory.

The current proposed solution for determining workgroup sizes:

1. Get work group size from the autodiscovery
2. get pseudochannel size from auto discovery

Given the board <pseudochannel_size> and <workgroup_size> from autodiscovery or a constant in runtime (as long it is determined during compile time).
And each work item writes that is <type_size> bytes vectorized by
Then each workgroup's size is <wg_space> = <workgroup_size> * <type_size> * .
need space for input and output buffers which two separate buffers both within one global memory address space. So size of 1 of these buffers is therefore <pseudochannel_size> / 2.

Then the formula for determining number of work groups constrained by size of global memory is: <num_workgroup> = (<pseudochannel_size> / 2) / <wg_space>

[sherry-yuan]

Thanks Artem @artemrad for the info! There is no need for clEnqueueNDRange. clEnqueueTask will do what we wanted (without the need to know workgroup size)

Precise answer below
"""
In general you do no need WG size. Launch this kernel as a task, rather than a NDRange. So default to {1, 1, 1} for WG size and WI sizes; More specifically do what you would do if a kernel was launched with clEnqueueTask() instead of clEnqueueNDRangeKernel()
"""

[sherry-yuan]

Push to 2022.4 given that's the overall new target. Next steps, run sycl l3 set on the device global change with autodiscovery change, merge in the autodiscovery+runtime change. May depend on pushing specs in.