Align OpenCL kernel arg table with SPIR-V

The OpenCL table describing legal kernel arguments had the opposite
order from the corresponding SPIR-V table.  The OpenCL table listed
the SYCL argument type in the left column while the SPIR-V table listed
the SYCL argument tupe in the right column.  I like the SPIR-V format
better, so change the OpenCL table to match.  This also adopts some of
the same wording from the SPIR-V table into the OpenCL one.

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc

@@ -149,46 +149,54 @@ sycl::kernel k = kb.ext_oneapi_get_kernel("foo");
 
 === Kernel argument restrictions
 
-When a kernel is defined in OpenCL C and invoked from SYCL via a `kernel`
-object, the arguments to the kernel are restricted to certain types.
-In general, the host application passes an argument value via
-`handler::set_arg` using one type and the kernel receives the argument value
-as a corresponding OpenCL C type.
-The following table lists the set of valid types for these kernel arguments:
-
+The following table defines the set of OpenCL C kernel argument types that are
+supported by this extension and explains how to pass each type of argument from
+SYCL.
 
 [%header,cols="1,1"]
 |===
-|Type in SYCL host code
-|Type in OpenCL C kernel
+|OpenCL C type
+|Corresponding SYCL type
 
-|One of the OpenCL scalar types (e.g. `cl_int`, `cl_float`, etc.)
-|The corresponding OpenCL C type (e.g. `int`, `float`, etc.)
+|One of the OpenCL C scalar types (e.g. `int`, `float`, etc.)
+|A {cpp} type that is device copyable, which has the same width and data
+ representation.
 
-|A USM pointer.
-|A `+__global+` pointer of the corresponding type.
+[_Note:_ Applications typically use the corresponding OpenCL type (e.g.
+`cl_int`, `cl_float`, etc.)
+_{endnote}_]
 
-|A class (or struct) that is device copyable in SYCL whose elements are
- composed of OpenCL scalar types or USM pointers.
-|A class (or struct) passed by value whose elements have the corresponding
- OpenCL C types.
+|A `+__global+` pointer.
+|Either a {cpp} pointer (typically a pointer to USM memory) or an `accessor`
+ whose target is `target::device`.
 
-|An `accessor` with `target::device` whose `DataT` is an OpenCL scalar type,
- a USM pointer, or a device copyable class (or struct) whose elements are
- composed of these types.
-|A `+__global+` pointer to the first element of the accessor's buffer.
- The pointer has the corresponding OpenCL C type.
+|A `+__local+` pointer.
+|A `local_accessor`.
 
-[_Note:_ The accessor's size is not passed as a kernel argument, so the host
-code must pass a separate argument with the size if this is desired.
+[_Note:_ The `local_accessor` merely conveys the size of the local memory, such
+that the kernel argument points to a local memory buffer of _N_ bytes, where
+_N_ is the value returned by `local_accessor::byte_size`.
+If the application wants to pass other information from the `local_accessor` to
+the kernel (such as the value _N_), it must pass this as separate kernel
+arguments.
 _{endnote}_]
 
-|A `local_accessor` whose `DataT` is an OpenCL scalar type, a USM pointer, or a
- device copyable class (or struct) whose elements are composed of these types.
-|A `+__local+` pointer to the first element of the accessor's local memory.
- The pointer has the corresponding OpenCL C type.
+|A class (or struct) passed by value.
+|A {cpp} struct or class that is device copyable, which has the same size and
+ data representation as the OpenCL C struct.
+
+[_Note:_ The SYCL argument must not contain any `accessor` or `local_accessor`
+members because these types are not device copyable.
+If the OpenCL C structure contains a pointer member, the corresponding SYCL
+structure member is typically a USM pointer.
+_{endnote}_]
 |===
 
+When data allocated on the host is accessed by the kernel via a pointer, the
+application must ensure that the data has the same size and representation on
+the host and inside the OpenCL C kernel.
+Applications can use the OpenCL types (e.g. `cl_int`) for this purpose.
+
 
 == Example