AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
77bb1c5616b8d1cd2bf3b22cb9f251f960b237ba
rocm-systems/opencl/api/opencl/amdocl/cl_execute.cpp
T
foreman 77bb1c5616 initial commit
2014-07-04 16:17:05 -04:00

1044 lines
41 KiB
C++
Raw Blame History

//
// Copyright (c) 2008 Advanced Micro Devices, Inc. All rights reserved.
//
#include "cl_common.hpp"
#include "platform/kernel.hpp"
#include "platform/ndrange.hpp"
#include "platform/command.hpp"
#include "platform/program.hpp"
#include <icd/icd_dispatch.h>
/*! \addtogroup API
* @{
*
* \addtogroup CL_Exec Executing Kernel Objects
*
* @{
*/
/*! \brief Enqueue a command to execute a kernel on a device.
*
* \param command_queue is a valid command-queue. The kernel will be queued
* for execution on the device associated with \a command_queue.
*
* \param kernel is a valid kernel object. The OpenCL context associated with
* \a kernel and \a command-queue must be the same.
*
* \param work_dim is the number of dimensions used to specify the global
* work-items and work-items in the work-group. \a work_dim must be greater
* than zero and less than or equal to three.
*
* \param global_work_offset must currently be a NULL value. In a future
* revision of OpenCL, \a global_work_offset can be used to specify an array
* of \a work_dim unsigned values that describe the offset used to calculate
* the global ID of a work-item instead of having the global IDs always start
* at offset (0, 0, 0).
*
* \param global_work_size points to an array of \a work_dim unsigned values
* that describe the number of global work-items in \a work_dim dimensions
* that will execute the kernel function. The total number of global
* work-items is computed as global_work_size[0] * ...
* * global_work_size[work_dim - 1].
*
* \param local_work_size points to an array of \a work_dim unsigned values
* that describe the number of work-items that make up a work-group (also
* referred to as the size of the work-group) that will execue the kernel
* specified by kernel.
*
* \param num_events_in_wait_list specifies the number of event objects in
* \a event_wait_list
*
* \param event_wait_list specifies events that need to complete before this
* particular command can be executed. If \a event_wait_list is NULL, then
* this particular command does not wait on any event to complete.
* If \a event_wait_list is NULL, \a num_events_in_wait_list must be 0.
* If \a event_wait_list is not NULL, the list of events pointed to by
* \a event_wait_list must be valid and \a num_events_in_wait_list must be
* greater than 0. The events specified in \a event_wait_list act as
* synchronization points.
*
* \param event returns an event object that identifies this particular kernel
* execution instance. Event objects are unique and can be used to identify a
* particular kernel execution instance later on. If \a event is NULL, no
* event will be created for this kernel execution instance and therefore it
* will not be possible for the application to query or queue a wait for this
* particular kernel execution instance.
*
* The total number of work-items in a work-group is computed as
* local_work_size[0] * ... * local_work_size[work_dim - 1].
* The total number of work-items in the work-group must be less than or equal
* to the CL_DEVICE_MAX_WORK_GROUP_SIZE. The explicitly specified
* \a local_work_size will be used to determine how to break the global work-
* items specified by global_work_size into appropriate work-group instances.
* If \a local_work_size is specified, the values specified in
* \a global_work_size[0], ..., global_work_size[work_dim - 1] must be evenly
* divisable by the corresponding values specified in \a local_work_size[0],
* ..., local_work_size[work_dim - 1]. \a local_work_size can also be a NULL
* value in which case the OpenCL implementation will determine how to be
* break the global work-items into appropriate work-groups.
*
* If \a local_work_size is NULL and no work-group size is specified when the
* kernel is compiled, the OpenCL implementation will determine how to break
* the global work-items specified by \a global_work_size into appropriate
* work-group instances. The work-group size to be used for kernel can also be
* specified in the program source using the
* __attribute__((reqd_work_group_size(X, Y, Z))) qualifier. In this case the
* size of work group specified by \a local_work_size must match the value
* specified by the \a reqd_work_group_size attribute qualifier.
*
* These work-group instances are executed in parallel across multiple
* compute units or concurrently on the same compute unit. Each work-item
* is uniquely identified by a global identifier. The global ID, which can be
* read inside the kernel is computed using the value given by
* \a global_work_size and \a global_work_offset.
*
* \return One of the following values:
*
* - CL_SUCCESS if the kernel execution was successfully queued
*
* - CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built program
* executable available for device associated with \a command_queue.
*
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue
*
* - CL_INVALID_KERNEL if \a kernel is not a valid kernel object.
*
* - CL_INVALID_CONTEXT if context associated with command_queue and kernel are
* not the same or if the context associated with command_queue and events in
* event_wait_list are not the same.
*
* - CL_INVALID_KERNEL_ARGS if the kernel argument values have not been
* specified or are not valid for the device on which kernel will be
* executed.
*
* - CL_INVALID_WORK_DIMENSION if \a work_dim is not a valid value
* (i.e. a value between 1 and 3).
*
* - CL_INVALID_WORK_GROUP_SIZE if \a local_work_size is specified and number
* of workitems specified by \a global_work_size is not evenly divisable by
* size of work-given by \a local_work_size or does not match the work-group
* size specified for kernel using the
* __attribute__((reqd_work_group_size(X, Y, Z))) qualifier in program
* source.
*
* - CL_INVALID_GLOBAL_OFFSET if \a global_work_offset is not NULL.
*
* - CL_OUT_OF_RESOURCES if there is a failure to queue the execution instance
* of \a kernel on the command-queue because of insufficient resources
* needed to execute the kernel. For example, the explicitly specified
* \a local_work_dim in range causes a failure to execute the kernel because
* of insufficient resources such as registers or local memory. Another
* example would be the number of read-only image args used in kernel exceed
* the CL_DEVICE_MAX_READ_IMAGE_ARGS value for device or the number of
* write-only image args used in kernel exceed the
* CL_DEVICE_MAX_WRITE_IMAGE_ARGS value for device or the number of samplers
* used in kernel exceed CL_DEVICE_MAX_SAMPLERS for device.
*
* - CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate memory
* for image or buffer objects specified as arguments to kernel.
*
* - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
* \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
* \a num_events_in_wait_list is 0, or if event objects in
* \a event_wait_list are not valid events.
*
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources
* required by the runtime.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clEnqueueNDRangeKernel, (
cl_command_queue command_queue,
cl_kernel kernel,
cl_uint work_dim,
const size_t *global_work_offset,
const size_t *global_work_size,
const size_t *local_work_size,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event))
{
*not_null(event) = NULL;
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid(kernel)) {
return CL_INVALID_KERNEL;
}
amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
if (NULL == queue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue& hostQueue = *queue;
const amd::Kernel* amdKernel = as_amd(kernel);
if (&hostQueue.context() != &amdKernel->program().context()) {
return CL_INVALID_CONTEXT;
}
const amd::Device& device = hostQueue.device();
const device::Kernel* devKernel = amdKernel->getDeviceKernel(device);
if (devKernel == NULL) {
return CL_INVALID_PROGRAM_EXECUTABLE;
}
if (amdKernel->parameters().getSvmSystemPointersSupport() == FGS_YES &&
!(device.info().svmCapabilities_ & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM)) {
// The user indicated that this kernel will access SVM system pointers,
// but the device does not support them.
return CL_INVALID_OPERATION;
}
if (work_dim < 1 || work_dim > 3) {
return CL_INVALID_WORK_DIMENSION;
}
#if !defined(CL_VERSION_1_1)
if (global_work_offset != NULL) {
return CL_INVALID_GLOBAL_OFFSET;
}
#endif // CL_VERSION
if (global_work_size == NULL) {
return CL_INVALID_VALUE;
}
if (local_work_size == NULL) {
static size_t zeroes[3] = { 0, 0, 0 };
local_work_size = zeroes;
}
else {
size_t numWorkItems = 1;
for (cl_uint dim = 0; dim < work_dim; ++dim) {
if (local_work_size[dim] == 0 || local_work_size[dim]
> device.info().maxWorkItemSizes_[dim]) {
return CL_INVALID_WORK_ITEM_SIZE;
}
if ((local_work_size[dim] != 0) &&
(devKernel->workGroupInfo()->compileSize_[0] != 0) && (local_work_size[dim] !=
devKernel->workGroupInfo()->compileSize_[dim])) {
return CL_INVALID_WORK_GROUP_SIZE;
}
if ((global_work_size[dim] == 0) ||
(((global_work_size[dim] % local_work_size[dim]) != 0) &&
(!device.settings().partialDispatch_ ||
devKernel->workGroupInfo()->uniformWorkGroupSize_))) {
return CL_INVALID_WORK_GROUP_SIZE;
}
numWorkItems *= local_work_size[dim];
}
if (numWorkItems > devKernel->workGroupInfo()->size_) {
return CL_INVALID_WORK_GROUP_SIZE;
}
}
// Check that all parameters have been defined.
if (!amdKernel->parameters().check()) {
return CL_INVALID_KERNEL_ARGS;
}
// Check that we do not exceed the amount of available local memory.
const size_t align = device.info().minDataTypeAlignSize_;
cl_ulong requiredLocalMemSize =
static_cast<cl_ulong>(amdKernel->parameters().localMemSize(align)) +
amd::alignUp(devKernel->workGroupInfo()->localMemSize_, align);
if (requiredLocalMemSize > device.info().localMemSize_) {
return CL_OUT_OF_RESOURCES;
}
amd::Command::EventWaitList eventWaitList;
cl_int err = amd::clSetEventWaitList(eventWaitList,
hostQueue.context(), num_events_in_wait_list, event_wait_list);
if (err != CL_SUCCESS) {
return err;
}
amd::NDRangeContainer ndrange((size_t) work_dim,
global_work_offset, global_work_size, local_work_size);
amd::NDRangeKernelCommand* command = new amd::NDRangeKernelCommand(
hostQueue, eventWaitList, *as_amd(kernel), ndrange);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
// ndrange is now owned by command. Do not delete it!
// Make sure we have memory for the command execution
cl_int result = command->validateMemory();
if (result != CL_SUCCESS) {
delete command;
return result;
}
command->enqueue();
*not_null(event) = as_cl(&command->event());
if (event == NULL) {
command->release();
}
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! \brief Enqueue a command to execute a kernel on a device.
* The kernel is executed using a single work-item.
*
* \param command_queue is a valid command-queue. The kernel will be queued
* for execution on the device associated with \a command_queue.
*
* \param kernel is a valid kernel object. The OpenCL context associated with
* \a kernel and \a command-queue must be the same.
*
* \param num_events_in_wait_list specifies the number of event objects in
* \a event_wait_list
*
* \param event_wait_list specifies events that need to complete before this
* particular command can be executed. If \a event_wait_list is NULL, then
* this particular command does not wait on any event to complete.
* If \a event_wait_list is NULL, \a num_events_in_wait_list must be 0.
* If \a event_wait_list is not NULL, the list of events pointed to by
* \a event_wait_list must be valid and \a num_events_in_wait_list must be
* greater than 0. The events specified in \a event_wait_list act as
* synchronization points.
*
* \param event returns an event objects that identifies this particular kernel
* execution instance. Event objects are unique and can be used to identify a
* particular kernel execution instance later on. If \a event is NULL, no event
* will be created for this kernel execution instance and therefore it will not
* be possible for the application to query or queue a wait for this particular
* kernel execution instance.
*
* \return One of the following values:
* - CL_SUCCESS if the kernel execution was successfully queued.
* - CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built program
* executable available for device associated with \a command_queue.
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue.
* - CL_INVALID_KERNEL if \a kernel is not a valid kernel object.
* - CL_INVALID_KERNEL_ARGS if the kernel argument values have not been
* specified or are not valid for the device on which kernel will be
* executed.
* - CL_INVALID_WORK_GROUP_SIZE if a work-group size is specified for
* kernel using the __attribute__((reqd_work_group_size(X, Y, Z)))
* qualifier in program source and is not (1, 1, 1).
* - CL_OUT_OF_RESOURCES if there is a failure to queue the execution instance
* of kernel on the command-queue because of insufficient resources needed
* to execute the kernel. For example, the explicitly specified
* \a local_work_dim in range causes a failure to execute the kernel because
* of insufficient resources such as registers or local memory. Another
* example would be the number of read-only image args used in kernel exceed
* the CL_DEVICE_MAX_READ_IMAGE_ARGS value for device or the number of
* write-only image args used in kernel exceed the
* CL_DEVICE_MAX_WRITE_IMAGE_ARGS value for device or the number of samplers
* used in kernel exceed CL_DEVICE_MAX_SAMPLERS for device.
* - CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate memory
* for image or buffer objects specified as arguments to kernel.
* - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
* \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
* \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
* are not valid events.
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
* by the runtime.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clEnqueueTask, (
cl_command_queue command_queue,
cl_kernel kernel,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event))
{
static size_t const globalWorkSize[3] = {1, 0, 0};
static size_t const localWorkSize[3] = {1, 0, 0};
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* hostQueue = as_amd(command_queue)->asHostQueue();
if (NULL == hostQueue) {
return CL_INVALID_COMMAND_QUEUE;
}
return hostQueue->dispatch_->clEnqueueNDRangeKernel(
command_queue, kernel, 1, NULL, globalWorkSize, localWorkSize,
num_events_in_wait_list, event_wait_list, event);
}
RUNTIME_EXIT
/*! \brief Enqueue a command to execute a native C/C++ function not compiled
* using the OpenCL compiler.
*
* \param command_queue is a valid command-queue. A native user function can
* only be executed on a command-queue created on a device that has
* CL_EXEC_NATIVE_KERNEL capability set in CL_DEVICE_EXECUTION_CAPABILITIES.
*
* \param user_func is a pointer to a host-callable user function.
*
* \param args is a pointer to the args list that \a user_func should be called
* with.
*
* \param cb_args is the size in bytes of the args list that args points to.
* The data pointed to by \a args and \a cb_args bytes in size will be copied
* and a pointer to this copied region will be passed to \a user_func. The copy
* needs to be done because the memory objects (cl_mem values) that args may
* contain need to be modified and replaced by appropriate pointers to global
* memory. When clEnqueueNativeKernel returns, the memory region pointed to by
* args can be reused by the application.
*
* \param num_mem_objects is the number of buffer objects that are passed in
* args.
*
* \param mem_list is a list of valid buffer objects, if \a num_mem_objects > 0
*
* \param args_mem_loc is a pointer to appropriate locations that args points
* to where memory object handles (cl_mem values) are stored. Before the user
* function is executed, the memory object handles are replaced by pointers to
* global memory.
*
* \param num_events_in_wait_list specifies the number of event objects in
* \a event_wait_list
*
* \param event_wait_list as described in clEnqueueNDRangeKernel.
*
* \param event returns an event objects that identifies this particular kernel
* execution instance. Event objects are unique and can be used to identify a
* particular kernel execution instance later on. If \a event is NULL, no event
* will be created for this kernel execution instance and therefore it will not
* be possible for the application to query or queue a wait for this particular
* kernel execution instance.
*
* \return One of the following values:
* - CL_SUCCESS if the user function execution instance was successfully queued
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue.
* - CL_INVALID_VALUE if \a user_func is NULL, or if \a args is a NULL value
* and \a num_mem_objects > 0 or if \a num_mem_objects > 0 and \a mem_list
* is NULL.
* - CL_INVALID_OPERATION if device cannot execute the native kernel.
* - CL_INVALID_MEM_OBJECT if one or more memory objects specified in
* \a mem_list are not valid or are not buffer objects.
* - CL_OUT_OF_RESOURCES if there is a failure to queue the execution instance
* of kernel on the command-queue because of insufficient resources needed
* to execute the kernel.
* - CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate memory
* for buffer objects specified as arguments to \a kernel.
* - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
* \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
* \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
* are not valid events.
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
* by the runtime.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clEnqueueNativeKernel, (
cl_command_queue command_queue,
void (CL_CALLBACK * user_func)(void *),
void *args,
size_t cb_args,
cl_uint num_mem_objects,
const cl_mem *mem_list,
const void **args_mem_loc,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event))
{
*not_null(event) = NULL;
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
if (NULL == queue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue& hostQueue = *queue;
const amd::Device& device = hostQueue.device();
if (!(device.info().executionCapabilities_ & CL_EXEC_NATIVE_KERNEL)) {
return CL_INVALID_OPERATION;
}
if (user_func == NULL
|| (num_mem_objects > 0 && (mem_list == NULL || args_mem_loc == NULL))
|| (num_mem_objects == 0 && (mem_list != NULL || args_mem_loc != NULL))
|| (args == NULL && (cb_args > 0 || num_mem_objects > 0))
|| (args != NULL && cb_args == 0)) {
return CL_INVALID_VALUE;
}
amd::Command::EventWaitList eventWaitList;
cl_int err = amd::clSetEventWaitList(eventWaitList,
hostQueue.context(), num_events_in_wait_list, event_wait_list);
if (err != CL_SUCCESS){
return err;
}
for (size_t i = 0; i < num_mem_objects; ++i) {
cl_mem obj = mem_list[i];
if (!is_valid(obj)) {
return CL_INVALID_MEM_OBJECT;
}
}
amd::NativeFnCommand* command = new amd::NativeFnCommand(
hostQueue, eventWaitList,
user_func, args, cb_args, num_mem_objects, mem_list, args_mem_loc);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
command->enqueue();
*not_null(event) = as_cl(&command->event());
if (event == NULL) {
command->release();
}
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! @}
*
* \addtogroup CL_Order Out of order Execution of Kernels and Memory Commands
*
* The OpenCL functions that are submitted to a command-queue are queued in
* the order the calls are made but can be configured to execute in-order or
* out-of-order. The properties argument in clCreateCommandQueue can be used
* to specify the execution order.
*
* If the CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE property of a command-queue
* is not set, the commands queued to a command-queue execute in order.
* For example, if an application calls clEnqueueNDRangeKernel to execute
* kernel A followed by a clEnqueueNDRangeKernel to execute kernel B,
* the application can assume that kernel A finishes first and then kernel B
* is executed. If the memory objects output by kernel A are inputs to kernel B
* then kernel B will see the correct data in memory objects produced
* by execution of kernel A. If the CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
* property of a commandqueue is set, then there is no guarantee that kernel A
* will finish before kernel B starts execution.
*
* Applications can configure the commands queued to a command-queue to
* execute out-of-order by setting the CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
* property of the commandqueue. This can be specified when the command-queue
* is created or can be changed dynamically using clSetCommandQueueProperty.
* In out-of-order execution mode there is no guarantee that the queued
* commands will finish execution in the order they were queued. As there is
* no guarantee that kernels will be executed in order i.e. based on when
* the clEnqueueNDRangeKernel calls are made within a command-queue, it is
* therefore possible that an earlier clEnqueueNDRangeKernel call to execute
* kernel A identified by event A may execute and/or finish later than a
* clEnqueueNDRangeKernel call to execute kernel B which was called by the
* application at a later point in time. To guarantee a specific order of
* execution of kernels, a wait on a particular event (in this case event A)
* can be used. The wait for event A can be specified in the event_wait_list
* argument to clEnqueueNDRangeKernel for kernel B.
*
* In addition, a wait for events or a barrier function can be queued to the
* command-queue. The wait for events command ensures that previously queued
* commands identified by the list of events to wait for have finished before
* the next batch of commands is executed. The barrier ensures that all
* previously queued commands in a command-queue have finished execution
* before the next batch of commands is executed.
*
* Similarly, commands to read, write, copy or map memory objects that are
* queued after clEnqueueNDRangeKernel, clEnqueueTask or clEnqueueNativeKernel
* commands are not guaranteed to wait for kernels scheduled for execution
* to have completed (if the CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE property
* is set). To ensure correct ordering of commands, the event object returned
* by clEnqueueNDRangeKernel, clEnqueueTask or clEnqueueNativeKernel can be
* used to queue a wait for event or a barrier command can be queued that must
* complete before reads or writes to the memory object(s) occur.
*
* @{
*/
/*! \brief Enqueue a marker command to \a command_queue.
*
* The marker command returns an event which can be used by to queue a wait on
* this marker event i.e. wait for all commands queued before the marker
* command to complete.
*
* \return One of the following values:
* - CL_SUCCESS if the function is successfully executed
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue
* - CL_INVALID_VALUE if \a event is a NULL value
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
* by the runtime.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clEnqueueMarker, (
cl_command_queue command_queue,
cl_event *event))
{
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* hostQueue = as_amd(command_queue)->asHostQueue();
if (NULL == hostQueue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::Command* command = new amd::Marker(*hostQueue, true);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
command->enqueue();
*not_null(event) = as_cl(&command->event());
if (event == NULL) {
command->release();
}
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! \brief enqueues a marker command which waits for either a list of events
* to complete, or if the list is empty it waits for all commands previously
* enqueued in \a command_queue to complete before it completes. This command
* returns an event which can be waited on, i.e. this event can be waited on
* to insure that all events either in the \a event_wait_list or all
* previously enqueued commands, queued before this command to
* \a command_queue, have completed.
*
* \param command_queue is a valid command-queue.
*
* \param num_events_in_wait_list specifies the number of events given
* by \a event_wait_list.
*
* \param event_wait_list specifies events that need to complete before this
* particular command can be executed.
* If \a event_wait_list is NULL, \a num_events_in_wait_list must be 0.
* If \a event_wait_list is not NULL, the list of events pointed to by
* \a event_wait_list must be valid and \a num_events_in_wait_list must
* be greater than 0. The events specified in event_wait_list act as
* synchronization points. The context associated with events in
* \a event_wait_list and \a command_queue must be the same. The
* memory associated with \a event_wait_list can be reused or freed after
* the function returns.
* If \a event_wait_list is NULL, then this particular command waits until
* all previous enqueued commands to \a command_queue have completed.
*
* \param event returns an event object that identifies this particular
* kernel execution instance. Event objects are unique and can be used to
* identify this marker command later on.
*
* \return CL_SUCCESS if the function is successfully executed.
* Otherwise, it returns one of the following errors:
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid \a command-queue.
* - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
* \a num_events_in_wait_list > 0, or event_wait_list is not NULL and
* \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
* are not valid events.
* - 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.
*
* \version 1.2r07
*/
RUNTIME_ENTRY(cl_int, clEnqueueMarkerWithWaitList, (
cl_command_queue command_queue,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event))
{
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* hostQueue = as_amd(command_queue)->asHostQueue();
if (NULL == hostQueue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::Command::EventWaitList eventWaitList;
cl_int err = amd::clSetEventWaitList(eventWaitList,
hostQueue->context(), num_events_in_wait_list, event_wait_list);
if (err != CL_SUCCESS) {
return err;
}
amd::Command* command = new amd::Marker(*hostQueue, true, eventWaitList);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
command->enqueue();
*not_null(event) = as_cl(&command->event());
if (event == NULL) {
command->release();
}
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! \brief Enqueue a wait for a specific event or a list of events to complete
* before any future commands queued in the command-queue are executed.
*
* \param command_queue is a valid command-queue.
*
* \param num_events specifies the number of events given by \a event_list.
*
* \param event_list is the list of events. Each event in \a event_list must
* be a valid event object returned by a previous call to:
* - clEnqueueNDRangeKernel
* - clEnqueueTask
* - clEnqueueNativeKernel
* - clEnqueue{Read|Write|Map}{Buffer|Image}
* - clEnqueueCopy{Buffer|Image}
* - clEnqueueCopyBufferToImage
* - clEnqueueCopyImageToBuffer
* - clEnqueueMarker.
* The events specified in \a event_list act as synchronization points.
*
* \return One of the following values:
* - CL_SUCCESS if the function was successfully executed.
* - CL_INVALID_COMMAND_QUEUE if c\a ommand_queue is not a valid command-queue
* - CL_INVALID_VALUE if \a num_events is zero or \a event_list is NULL
* - CL_INVALID_EVENT if event objects specified in \a event_list are not valid
* events
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
* by the runtime.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clEnqueueWaitForEvents, (
cl_command_queue command_queue,
cl_uint num_events,
const cl_event *event_list))
{
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
if (NULL == queue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue& hostQueue = *queue;
amd::Command::EventWaitList eventWaitList;
cl_int err = amd::clSetEventWaitList(eventWaitList,
hostQueue.context(), num_events, event_list);
if (err != CL_SUCCESS){
return err;
}
amd::Command* command = new amd::Marker(hostQueue, false, eventWaitList);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
command->enqueue();
command->release();
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! \brief Enqueue a barrier operation.
*
* The clEnqueueBarrier command ensures that all queued commands in
* \a command_queue have finished execution before the next batch of commands
* can begin execution. clEnqueueBarrier is a synchronization point.
*
* \return One of the following values:
* - CL_SUCCESS if the function was executed successfully
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
* by the runtime.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clEnqueueBarrier, (cl_command_queue command_queue))
{
//! @todo: Unimplemented();
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! \brief enqueues a barrier command which waits for either a list of events
* to complete, or if the list is empty it waits for all commands previously
* enqueued in \a command_queue to complete before it completes. This command
* blocks command execution, that is, any following commands enqueued after it
* do not execute until it completes. This command returns an event which can
* be waited on, i.e. this event can be waited on to insure that all events
* either in the \a event_wait_list or all previously enqueued commands,
* queued before this command to command_queue, have completed
*
* \param command_queue is a valid command-queue.
*
* \param num_events_in_wait_list specifies the number of events given
* by \a event_wait_list.
*
* \param event_wait_list specifies events that need to complete before this
* particular command can be executed.
* If \a event_wait_list is NULL, \a num_events_in_wait_list must be 0.
* If \a event_wait_list is not NULL, the list of events pointed to by
* \a event_wait_list must be valid and \a num_events_in_wait_list must
* be greater than 0. The events specified in event_wait_list act as
* synchronization points. The context associated with events in
* \a event_wait_list and \a command_queue must be the same. The
* memory associated with \a event_wait_list can be reused or freed after
* the function returns.
* If \a event_wait_list is NULL, then this particular command waits until
* all previous enqueued commands to \a command_queue have completed.
*
* \param event returns an event object that identifies this particular
* kernel execution instance. Event objects are unique and can be used to
* identify this marker command later on.
*
* \return CL_SUCCESS if the function is successfully executed.
* Otherwise, it returns one of the following errors:
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid \a command-queue.
* - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
* \a num_events_in_wait_list > 0, or event_wait_list is not NULL and
* \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
* are not valid events.
* - 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.
*
* \version 1.2r07
*/
RUNTIME_ENTRY(cl_int, clEnqueueBarrierWithWaitList, (
cl_command_queue command_queue,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event))
{
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* hostQueue = as_amd(command_queue)->asHostQueue();
if (NULL == hostQueue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::Command::EventWaitList eventWaitList;
cl_int err = amd::clSetEventWaitList(eventWaitList,
hostQueue->context(), num_events_in_wait_list, event_wait_list);
if (err != CL_SUCCESS) {
return err;
}
//!@note: with the current runtime architecture and in-order execution
//! barrier and marker should be the same operation
amd::Command* command = new amd::Marker(*hostQueue, true, eventWaitList);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
command->enqueue();
*not_null(event) = as_cl(&command->event());
if (event == NULL) {
command->release();
}
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! @}
*
* \addtogroup CL_Profiling Profiling Operations on Memory Objects and Kernels
*
* Profiling of OpenCL functions that are enqueued as commands to a
* command-queue. The specific functions being referred to are:
* - clEnqueue{Read|Write|Map}Buffer,
* - clEnqueue{Read|Write|Map}Image,
* - clEnqueueCopy{Buffer|Image},
* - clEnqueueCopyImageToBuffer,
* - clEnqueueCopyBufferToImage,
* - clEnqueueNDRangeKernel ,
* - clEnqueueTask and
* - clEnqueueNativeKernel.
* These enqueued commands are identified by unique event objects.
*
* Event objects can be used to capture profiling information that measure
* execution time of a command. Profiling of OpenCL commands can be enabled
* either by using a command-queue created with CL_QUEUE_PROFILING_ENABLE
* flag set in properties arguments to clCreateCommandQueue or by setting the
* CL_QUEUE_PROFILING_ENABLE flag in properties arguments to
* clSetCommandQueueProperty.
*
* @{
*/
/*! \brief Return profiling information for the command associated with event.
*
* \param event specifies the event object.
*
* \param param_name specifies the profiling data to query.
*
* \param param_value is a pointer to memory where the appropriate result being
* queried is returned. If \a param_value is NULL, it is ignored.
*
* \param param_value_size is used to specify the size in bytes of memory
* pointed to by \a param_value. This size must be >= size of return type
*
* \param param_value_size_ret returns the actual size in bytes of data copied
* to \a param_value. If \a param_value_size_ret is NULL, it is ignored.
*
* The unsigned 64-bit values returned can be used to measure the time in
* nano-seconds consumed by OpenCL commands. OpenCL devices are required to
* correctly track time across changes in frequency and p-states. The
* CL_DEVICE_PROFILING_TIMER_RESOLUTION specifies the resolution of the timer
* i.e. the number of nanoseconds elapsed before the timer is incremented.
*
* \return One of the following values:
* - CL_SUCCESS if the function is executed successfully and the profiling
* information has been recorded
* - CL_PROFILING_INFO_NOT_AVAILABLE if the profiling information is currently
* not available (because the command identified by event has not completed)
* - CL_INVALID_VALUE if \a param_name is not valid, or if size in bytes
* specified by param_value_size is < size of return type and \a param_value
* is not NULL
* - CL_INVALID_EVENT if \a event is a not a valid event object.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clGetEventProfilingInfo, (
cl_event event,
cl_profiling_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret))
{
if (!is_valid(event)) {
return CL_INVALID_EVENT;
}
if (!as_amd(event)->profilingInfo().enabled_) {
return CL_PROFILING_INFO_NOT_AVAILABLE;
}
if (param_value != NULL && param_value_size < sizeof(cl_ulong)) {
return CL_INVALID_VALUE;
}
*not_null(param_value_size_ret) = sizeof(cl_ulong);
if (param_value != NULL) {
cl_ulong value = 0;
switch (param_name) {
case CL_PROFILING_COMMAND_END:
value = as_amd(event)->profilingInfo().end_;
break;
case CL_PROFILING_COMMAND_START:
value = as_amd(event)->profilingInfo().start_;
break;
case CL_PROFILING_COMMAND_SUBMIT:
value = as_amd(event)->profilingInfo().submitted_;
break;
case CL_PROFILING_COMMAND_QUEUED:
value = as_amd(event)->profilingInfo().queued_;
break;
default:
return CL_INVALID_VALUE;
}
if (value == 0) {
return CL_PROFILING_INFO_NOT_AVAILABLE;
}
*(cl_ulong*)param_value = value;
}
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! @}
* \addtogroup CL_FlushFinish Flush and Finish
* @{
*/
/*! \brief Issue all previously queued OpenCL commands in \a command_queue to
* the device associated with command_queue.
*
* clFlush only guarantees that all queued commands to \a command_queue get
* issued to the appropriate device. There is no guarantee that they will be
* complete after clFlush returns.
*
* \return One of the following values:
* - CL_SUCCESS if the function call was executed successfully
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
* by the runtime.
*
* Any blocking commands queued in a command-queue such as
* clEnqueueRead{Image|Buffer} with \a blocking_read set to CL_TRUE,
* clEnqueueWrite{Image|Buffer} with \a blocking_write set to CL_TRUE,
* clEnqueueMap{Buffer|Image} with \a blocking_map set to CL_TRUE or
* clWaitForEvents perform an implicit flush of the command-queue.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clFlush, (cl_command_queue command_queue))
{
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* hostQueue = as_amd(command_queue)->asHostQueue();
if (NULL == hostQueue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::Command* command = new amd::Marker(*hostQueue, false);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
command->enqueue();
command->release();
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! \brief Block until all previously queued OpenCL runtime commands in
* \a command_queue are issued to the associated device and have completed.
*
* clFinish does not return until all queued commands in \a command_queue have
* been processed and completed. clFinish is also a synchronization point.
*
* \return One of the following values:
* - CL_SUCCESS if the function call was executed successfully.
* - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue
* - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
* by the runtime.
*
* \version 1.0r33
*/
RUNTIME_ENTRY(cl_int, clFinish, (cl_command_queue command_queue))
{
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue* hostQueue = as_amd(command_queue)->asHostQueue();
if (NULL == hostQueue) {
return CL_INVALID_COMMAND_QUEUE;
}
hostQueue->finish();
return CL_SUCCESS;
}
RUNTIME_EXIT
/*! @}
* @}
*/
Reference in New Issue View Git Blame Copy Permalink