From 97a2094fb2fff6d85674f90c364d6a896fa39c7d Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Thu, 29 Aug 2024 16:23:02 -0700 Subject: [PATCH 1/4] rephrase and correct the descriptions for clSetKernelExecInfo --- api/opencl_runtime_layer.asciidoc | 137 +++++++++++------------------- 1 file changed, 48 insertions(+), 89 deletions(-) diff --git a/api/opencl_runtime_layer.asciidoc b/api/opencl_runtime_layer.asciidoc index 438957c5..c4de581b 100644 --- a/api/opencl_runtime_layer.asciidoc +++ b/api/opencl_runtime_layer.asciidoc @@ -10433,16 +10433,16 @@ Otherwise, it returns one of the following errors: required by the OpenCL implementation on the host. -- -[open,refpage='clSetKernelExecInfo',desc='Pass additional information other than argument values to a kernel.',type='protos'] +[open,refpage='clSetKernelExecInfo',desc='Set additional information specifying how a kernel will execute.',type='protos'] -- -To pass additional information other than argument values to a kernel, call +To set additional information specifying how a kernel will execute, call the function include::{generated}/api/protos/clSetKernelExecInfo.txt[] include::{generated}/api/version-notes/clSetKernelExecInfo.asciidoc[] - * _kernel_ specifies the kernel object being queried. - * _param_name_ specifies the information to be passed to kernel. + * _kernel_ is a valid kernel object. + * _param_name_ specifies the type of information to set. The list of supported _param_name_ types and the corresponding values passed in _param_value_ is described in the <> table. @@ -10460,22 +10460,42 @@ include::{generated}/api/version-notes/clSetKernelExecInfo.asciidoc[] include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_PTRS.asciidoc[] | {void_TYPE}*[] - | SVM pointers must reference locations contained entirely within - buffers that are passed to kernel as arguments, or that are passed - through the execution information. - - Non-argument SVM buffers must be specified by passing pointers to - those buffers via {clSetKernelExecInfo} for coarse-grain and - fine-grain buffer SVM allocations but not for finegrain system SVM - allocations. + | Specifies a set of pointers to SVM allocations that may be accessed + by the kernel in addition to those set directly as kernel arguments. + Each of the pointers can be the pointer returned by {clSVMAlloc} or can + be a pointer offset into the SVM region. + It is sufficient to provide one pointer for each SVM allocation. + + The complete set of pointers is specified by each call to + {clSetKernelExecInfo} and replaces any previously specified set of + pointers. + To specify that no SVM allocations will be accessed by a kernel other + than those passed as kernel arguments, specify an empty set by passing + _param_value_size_ equal to zero and _param_value_ equal to `NULL`. + + Non-argument pointers to SVM allocations must be specified for + coarse-grain and fine-grain buffer SVM allocations, but not for + fine-grain system SVM allocations. | {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM_anchor} include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM.asciidoc[] | {cl_bool_TYPE} - | 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_. + | Specifies whether the kernel may use pointers to system allocations. + + When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_FALSE}, the + OpenCL implementation may assume that system pointers are not passed as + kernel arguments and are not stored inside SVM allocations passed as + kernel arguments. + + When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE}, the + OpenCL implementation must assume that system pointers might be passed + as kernel arguments or stored inside SVM allocations passed as + kernel arguments. + + If {clSetKernelExecInfo} has not been called with a value for + {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM}, the default value is + {CL_TRUE} if the device on which the kernel is enqueued supports + fine-grain system SVM, otherwise the default is {CL_FALSE}. |==== // refError @@ -10486,80 +10506,19 @@ Otherwise, it returns one of the following errors: * {CL_INVALID_KERNEL} if _kernel_ is a not a valid kernel object. * {CL_INVALID_OPERATION} if no devices in the context associated with _kernel_ support SVM. - * {CL_INVALID_VALUE} if _param_name_ is not valid, if _param_value_ is - `NULL` or if the size specified by _param_value_size_ is not valid. * {CL_INVALID_OPERATION} if _param_name_ is {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} and _param_value_ is {CL_TRUE} - but no devices in context associated with _kernel_ support fine-grain + and no devices in the context associated with _kernel_ support fine-grain system SVM allocations. + * {CL_INVALID_VALUE} if _param_name_ is not valid, if _param_value_ is + `NULL` and _param_value_size_ is greater than zero, or if the size specified + by _param_value_size_ is not valid. * {CL_OUT_OF_RESOURCES} if there is a failure to allocate resources required by the OpenCL implementation on the device. * {CL_OUT_OF_HOST_MEMORY} if there is a failure to allocate resources required by the OpenCL implementation on the host. -- -[NOTE] -==== -Coarse-grain or fine-grain buffer SVM pointers used by a kernel which -are not passed as a kernel arguments must be specified using -{clSetKernelExecInfo} with {CL_KERNEL_EXEC_INFO_SVM_PTRS}. -For example, if SVM buffer A contains a pointer to another SVM buffer B, -and the kernel dereferences that pointer, then a pointer to B must -either be passed as an argument in the call to that kernel or it must be -made available to the kernel using {clSetKernelExecInfo}. -For example, we might pass extra SVM pointers as follows: - -[source,opencl] ----- -clSetKernelExecInfo(kernel, - CL_KERNEL_EXEC_INFO_SVM_PTRS, - num_ptrs * sizeof(void *), - extra_svm_ptr_list); ----- - -Here `num_ptrs` specifies the number of additional SVM pointers while -`extra_svm_ptr_list` specifies a pointer to memory containing those SVM -pointers. - -When calling {clSetKernelExecInfo} with {CL_KERNEL_EXEC_INFO_SVM_PTRS} to -specify pointers to non-argument SVM buffers as extra arguments to a kernel, -each of these pointers can be the SVM pointer returned by {clSVMAlloc} or -can be a pointer + offset into the SVM region. -It is sufficient to provide one pointer for each SVM buffer used. - -{CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is used to indicate whether -SVM pointers used by a kernel will refer to system allocations or not. - -{CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} = {CL_FALSE} indicates that the -OpenCL implementation may assume that system pointers are not passed as -kernel arguments and are not stored inside SVM allocations passed as kernel -arguments. - -{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. -In this case, if the device to which the kernel is enqueued does not support -system SVM pointers, {clEnqueueNDRangeKernel} and {clEnqueueTask} will return a -{CL_INVALID_OPERATION} error. -If none of the devices in the context associated with kernel support -fine-grain system SVM allocations, {clSetKernelExecInfo} will return a -{CL_INVALID_OPERATION} error. - -If {clSetKernelExecInfo} has not been called with a value for -{CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM}, the default value is used for -this kernel attribute. -The default value depends on whether the device on which the kernel is -enqueued supports fine-grain system SVM allocations. -If so, the default value used is {CL_TRUE} (system pointers might be passed); -otherwise, the default is {CL_FALSE}. - -A call to {clSetKernelExecInfo} for a given value of _param_name_ -replaces any prior value passed for that value of _param_name_. -Only one _param_value_ will be stored for each value of _param_name_. -==== - - === Copying Kernel Objects NOTE: Copying kernel objects is <> version 2.1. @@ -11434,10 +11393,10 @@ Otherwise, it returns one of the following errors: _num_events_in_wait_list_ is 0, or if event objects in _event_wait_list_ are not valid events. * {CL_INVALID_OPERATION} if SVM pointers are passed as arguments to a kernel - and the device does not support SVM or if system pointers are passed as - arguments to a kernel and/or stored inside SVM allocations passed as - kernel arguments and the device does not support fine grain system SVM - allocations. + and the device does not support SVM, or if system pointers are passed as + arguments to a kernel and the device does not support fine-grain system SVM, + or if {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE} and the device + does not support fine-grain system SVM. * {CL_OUT_OF_RESOURCES} if there is a failure to allocate resources required by the OpenCL implementation on the device. * {CL_OUT_OF_HOST_MEMORY} if there is a failure to allocate resources @@ -11529,10 +11488,10 @@ Otherwise, it returns one of the following errors: _num_events_in_wait_list_ is 0, or if event objects in _event_wait_list_ are not valid events. * {CL_INVALID_OPERATION} if SVM pointers are passed as arguments to a kernel - and the device does not support SVM or if system pointers are passed as - arguments to a kernel and/or stored inside SVM allocations passed as - kernel arguments and the device does not support fine grain system SVM - allocations. + and the device does not support SVM, or if system pointers are passed as + arguments to a kernel and the device does not support fine-grain system SVM, + or if {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE} and the device + does not support fine-grain system SVM. * {CL_OUT_OF_RESOURCES} if there is a failure to allocate resources required by the OpenCL implementation on the device. * {CL_OUT_OF_HOST_MEMORY} if there is a failure to allocate resources From 1644e4a5216c0937edf03e19647da8a9543d212a Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Tue, 3 Sep 2024 22:25:37 -0700 Subject: [PATCH 2/4] further wordsmithing clarify that CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM does not affect kernel arguments --- api/opencl_runtime_layer.asciidoc | 30 ++++++++++++++++-------------- 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/api/opencl_runtime_layer.asciidoc b/api/opencl_runtime_layer.asciidoc index c4de581b..6e72ea3e 100644 --- a/api/opencl_runtime_layer.asciidoc +++ b/api/opencl_runtime_layer.asciidoc @@ -10433,10 +10433,9 @@ Otherwise, it returns one of the following errors: required by the OpenCL implementation on the host. -- -[open,refpage='clSetKernelExecInfo',desc='Set additional information specifying how a kernel will execute.',type='protos'] +[open,refpage='clSetKernelExecInfo',desc='Set additional execution information for a kernel.',type='protos'] -- -To set additional information specifying how a kernel will execute, call -the function +To set additional execution information for a kernel, call the function include::{generated}/api/protos/clSetKernelExecInfo.txt[] include::{generated}/api/version-notes/clSetKernelExecInfo.asciidoc[] @@ -10463,14 +10462,18 @@ include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_PTRS.asciidoc[] | Specifies a set of pointers to SVM allocations that may be accessed by the kernel in addition to those set directly as kernel arguments. Each of the pointers can be the pointer returned by {clSVMAlloc} or can - be a pointer offset into the SVM region. - It is sufficient to provide one pointer for each SVM allocation. + be a pointer to the middle of an SVM allocation. + It is sufficient to specify one pointer for each SVM allocation. + + Behavior is undefined if the kernel accesses a coarse-grain or + fine-grain buffer SVM allocation that is not set as a kernel argument + and is not in the set specified by {CL_KERNEL_EXEC_INFO_SVM_PTRS}. The complete set of pointers is specified by each call to {clSetKernelExecInfo} and replaces any previously specified set of pointers. To specify that no SVM allocations will be accessed by a kernel other - than those passed as kernel arguments, specify an empty set by passing + than those set as kernel arguments, specify an empty set by passing _param_value_size_ equal to zero and _param_value_ equal to `NULL`. Non-argument pointers to SVM allocations must be specified for @@ -10480,17 +10483,16 @@ include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_PTRS.asciidoc[] include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM.asciidoc[] | {cl_bool_TYPE} - | Specifies whether the kernel may use pointers to system allocations. + | Specifies whether the kernel may use pointers to system allocations + that are not set directly as kernel arguments. - When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_FALSE}, the - OpenCL implementation may assume that system pointers are not passed as - kernel arguments and are not stored inside SVM allocations passed as - kernel arguments. + When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_FALSE}, behavior + is undefined if the kernel acceses a system allocation that is not set as + a kernel argument. When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE}, the - OpenCL implementation must assume that system pointers might be passed - as kernel arguments or stored inside SVM allocations passed as - kernel arguments. + kernel may access system allocations that are not set directly as kernel + arguments. If {clSetKernelExecInfo} has not been called with a value for {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM}, the default value is From ee738f0c8d59e44ecdf3aa26188254126ee428f7 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 4 Sep 2024 10:50:27 -0700 Subject: [PATCH 3/4] fix typo --- api/opencl_runtime_layer.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/api/opencl_runtime_layer.asciidoc b/api/opencl_runtime_layer.asciidoc index 6e72ea3e..80497e31 100644 --- a/api/opencl_runtime_layer.asciidoc +++ b/api/opencl_runtime_layer.asciidoc @@ -10487,7 +10487,7 @@ include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM that are not set directly as kernel arguments. When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_FALSE}, behavior - is undefined if the kernel acceses a system allocation that is not set as + is undefined if the kernel accesses a system allocation that is not set as a kernel argument. When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE}, the From 8ce5495c275ade7e782992b9b7bb8fdf7aafd976 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 21 Oct 2024 18:15:55 -0700 Subject: [PATCH 4/4] simplify CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM --- api/opencl_runtime_layer.asciidoc | 27 ++++++++++++--------------- 1 file changed, 12 insertions(+), 15 deletions(-) diff --git a/api/opencl_runtime_layer.asciidoc b/api/opencl_runtime_layer.asciidoc index 80497e31..0484af19 100644 --- a/api/opencl_runtime_layer.asciidoc +++ b/api/opencl_runtime_layer.asciidoc @@ -10484,20 +10484,21 @@ include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_PTRS.asciidoc[] include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM.asciidoc[] | {cl_bool_TYPE} | Specifies whether the kernel may use pointers to system allocations - that are not set directly as kernel arguments. + that are not set directly as kernel arguments on devices that support + fine-grain system SVM allocations. - When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_FALSE}, behavior - is undefined if the kernel accesses a system allocation that is not set as - a kernel argument. + When a device supports fine-grain system SVM allocations and + {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE}, the kernel may + access system allocations that are not set directly as kernel arguments. - When {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE}, the - kernel may access system allocations that are not set directly as kernel - arguments. + Otherwise, if a device does not support fine-grain system SVM + allocations or when {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is + {CL_FALSE}, behavior is undefined if the kernel accesses a system + allocation that is not set as a kernel argument. If {clSetKernelExecInfo} has not been called with a value for {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM}, the default value is - {CL_TRUE} if the device on which the kernel is enqueued supports - fine-grain system SVM, otherwise the default is {CL_FALSE}. + {CL_TRUE}. |==== // refError @@ -11396,9 +11397,7 @@ Otherwise, it returns one of the following errors: are not valid events. * {CL_INVALID_OPERATION} if SVM pointers are passed as arguments to a kernel and the device does not support SVM, or if system pointers are passed as - arguments to a kernel and the device does not support fine-grain system SVM, - or if {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE} and the device - does not support fine-grain system SVM. + arguments to a kernel and the device does not support fine-grain system SVM. * {CL_OUT_OF_RESOURCES} if there is a failure to allocate resources required by the OpenCL implementation on the device. * {CL_OUT_OF_HOST_MEMORY} if there is a failure to allocate resources @@ -11491,9 +11490,7 @@ Otherwise, it returns one of the following errors: are not valid events. * {CL_INVALID_OPERATION} if SVM pointers are passed as arguments to a kernel and the device does not support SVM, or if system pointers are passed as - arguments to a kernel and the device does not support fine-grain system SVM, - or if {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} is {CL_TRUE} and the device - does not support fine-grain system SVM. + arguments to a kernel and the device does not support fine-grain system SVM. * {CL_OUT_OF_RESOURCES} if there is a failure to allocate resources required by the OpenCL implementation on the device. * {CL_OUT_OF_HOST_MEMORY} if there is a failure to allocate resources