does CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM also apply to kernel arguments?

[bashbaug]

Found while working on the fix for #1152.

The description for CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM is currently (emphasis mine):

This flag indicates whether the kernel uses pointers that are fine grain system SVM allocations. These fine grain system SVM pointers may be passed as arguments or defined in SVM buffers that are passed as arguments to kernel.

Additionally, in the "note" below:

CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM = CL_TRUE indicates that the OpenCL implementation must assume that system pointers might be passed as kernel arguments and/or stored inside SVM allocations passed as kernel arguments.

So, it seems like the setting for CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM should also apply to kernel arguments, but I don't see any description how the setting has any affect on or interaction with kernel arguments.

Here are a few specific questions (pseudocode):

// Assume that CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM for kernel is initially CL_TRUE:

// Allocate some memory.
void* ptr = malloc(...);

// Assume setting a system SVM kernel argument is initially successful:
clSetKernelArgSVMPointer(kernel, ptr);

// Now disable CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM:
clSetKernelExecInfo(kernel, CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM, CL_FALSE);

// Now is setting a system SVM kernel argument an error?  If so, which error?
// (see also note below)
clSetKernelArgSVMPointer(kernel, ptr);

// Is it OK to enqueue the kernel with a system SVM kernel argument or is this an error?
clEnqueueNDRangeKernel(kernel);

Note: the default value for CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM is also a little problematic, since the current specification says it "depends on whether the device on which the kernel is enqueued". If setting a system SVM kernel argument can be an error, and kernel was compiled for a multi-device context, and only one of the devices supports system SVM, we don't know which device the kernel is going be enqueued upon when we're setting kernel arguments...