Skip to content

A
angle
Commit 55df3ec1 by John Plate Committed by Angle LUCI CQ
Options

CL: Add remaining enqueue commands

Add support for remaining OpenCL 1.2 enqueue commands to front end and pass-through back end. Bug: angleproject:6015 Change-Id: Iab650e42d51e2105dc826088d3606c56d5cd0fd5 Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/2944966Reviewed-by: 's avatarJamie Madill <jmadill@chromium.org> Reviewed-by: 's avatarCody Northrop <cnorthrop@google.com> Commit-Queue: John Plate <jplate@google.com>
parent 5e5d17cd
Showing with 983 additions and 202 deletions
include/angle_cl.h
......@@ -33,6 +33,8 @@ using ProgramCB = void(CL_CALLBACK *)(cl_program program, void *user_data);
using EventCB = void(CL_CALLBACK *)(cl_event event, cl_int event_command_status, void *user_data);
using UserFunc = void(CL_CALLBACK *)(void *args);
template <typename T = void>
struct Dispatch
{
......
src/libANGLE/CLCommandQueue.cpp
......@@ -12,12 +12,32 @@
#include "libANGLE/CLDevice.h"
#include "libANGLE/CLEvent.h"
#include "libANGLE/CLImage.h"
#include "libANGLE/CLKernel.h"
#include "libANGLE/CLMemory.h"
#include <cstring>
namespace cl
{
namespace
{
void CheckCreateEvent(CommandQueue &queue,
cl_command_type commandType,
const rx::CLEventImpl::CreateFunc &createFunc,
cl_event *event,
cl_int &errorCode)
{
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(createFunc);
*event = Object::Create<Event>(errorCode, queue, commandType, createFunc);
}
}
} // namespace
cl_int CommandQueue::getInfo(CommandQueueInfo name,
size_t valueSize,
void *value,
......@@ -128,11 +148,7 @@ cl_int CommandQueue::enqueueReadBuffer(cl_mem buffer,
cl_int errorCode =
mImpl->enqueueReadBuffer(buf, blocking, offset, size, ptr, waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_READ_BUFFER, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_READ_BUFFER, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -155,11 +171,7 @@ cl_int CommandQueue::enqueueWriteBuffer(cl_mem buffer,
cl_int errorCode =
mImpl->enqueueWriteBuffer(buf, blocking, offset, size, ptr, waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_WRITE_BUFFER, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_WRITE_BUFFER, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -188,12 +200,7 @@ cl_int CommandQueue::enqueueReadBufferRect(cl_mem buffer,
buf, blocking, bufferOrigin, hostOrigin, region, bufferRowPitch, bufferSlicePitch,
hostRowPitch, hostSlicePitch, ptr, waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event =
Object::Create<Event>(errorCode, *this, CL_COMMAND_READ_BUFFER_RECT, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_READ_BUFFER_RECT, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -222,12 +229,7 @@ cl_int CommandQueue::enqueueWriteBufferRect(cl_mem buffer,
buf, blocking, bufferOrigin, hostOrigin, region, bufferRowPitch, bufferSlicePitch,
hostRowPitch, hostSlicePitch, ptr, waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event =
Object::Create<Event>(errorCode, *this, CL_COMMAND_WRITE_BUFFER_RECT, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_WRITE_BUFFER_RECT, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -250,11 +252,7 @@ cl_int CommandQueue::enqueueCopyBuffer(cl_mem srcBuffer,
cl_int errorCode = mImpl->enqueueCopyBuffer(src, dst, srcOffset, dstOffset, size, waitEvents,
eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_COPY_BUFFER, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_COPY_BUFFER, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -282,12 +280,7 @@ cl_int CommandQueue::enqueueCopyBufferRect(cl_mem srcBuffer,
srcRowPitch, srcSlicePitch, dstRowPitch,
dstSlicePitch, waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event =
Object::Create<Event>(errorCode, *this, CL_COMMAND_COPY_BUFFER_RECT, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_COPY_BUFFER_RECT, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -309,11 +302,7 @@ cl_int CommandQueue::enqueueFillBuffer(cl_mem buffer,
cl_int errorCode = mImpl->enqueueFillBuffer(buf, pattern, patternSize, offset, size, waitEvents,
eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_FILL_BUFFER, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_FILL_BUFFER, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -337,11 +326,7 @@ void *CommandQueue::enqueueMapBuffer(cl_mem buffer,
void *const map = mImpl->enqueueMapBuffer(buf, blocking, mapFlags, offset, size, waitEvents,
eventCreateFuncPtr, errorCode);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_MAP_BUFFER, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_MAP_BUFFER, eventCreateFunc, event, errorCode);
return map;
}
......@@ -366,11 +351,7 @@ cl_int CommandQueue::enqueueReadImage(cl_mem image,
cl_int errorCode = mImpl->enqueueReadImage(img, blocking, origin, region, rowPitch, slicePitch,
ptr, waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_READ_IMAGE, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_READ_IMAGE, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -396,11 +377,7 @@ cl_int CommandQueue::enqueueWriteImage(cl_mem image,
mImpl->enqueueWriteImage(img, blocking, origin, region, inputRowPitch, inputSlicePitch, ptr,
waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_WRITE_IMAGE, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_WRITE_IMAGE, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -423,11 +400,7 @@ cl_int CommandQueue::enqueueCopyImage(cl_mem srcImage,
cl_int errorCode = mImpl->enqueueCopyImage(src, dst, srcOrigin, dstOrigin, region, waitEvents,
eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_COPY_IMAGE, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_COPY_IMAGE, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -448,11 +421,7 @@ cl_int CommandQueue::enqueueFillImage(cl_mem image,
cl_int errorCode =
mImpl->enqueueFillImage(img, fillColor, origin, region, waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_FILL_IMAGE, eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_FILL_IMAGE, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -475,12 +444,7 @@ cl_int CommandQueue::enqueueCopyImageToBuffer(cl_mem srcImage,
cl_int errorCode = mImpl->enqueueCopyImageToBuffer(src, dst, srcOrigin, region, dstOffset,
waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_COPY_IMAGE_TO_BUFFER,
eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_COPY_IMAGE_TO_BUFFER, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -503,12 +467,7 @@ cl_int CommandQueue::enqueueCopyBufferToImage(cl_mem srcBuffer,
cl_int errorCode = mImpl->enqueueCopyBufferToImage(src, dst, srcOffset, dstOrigin, region,
waitEvents, eventCreateFuncPtr);
if (errorCode == CL_SUCCESS && event != nullptr)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_COPY_BUFFER_TO_IMAGE,
eventCreateFunc);
}
CheckCreateEvent(*this, CL_COMMAND_COPY_BUFFER_TO_IMAGE, eventCreateFunc, event, errorCode);
return errorCode;
}
......@@ -535,12 +494,203 @@ void *CommandQueue::enqueueMapImage(cl_mem image,
mImpl->enqueueMapImage(img, blocking, mapFlags, origin, region, imageRowPitch,
imageSlicePitch, waitEvents, eventCreateFuncPtr, errorCode);
if (errorCode == CL_SUCCESS && event != nullptr)
CheckCreateEvent(*this, CL_COMMAND_MAP_IMAGE, eventCreateFunc, event, errorCode);
return map;
}
cl_int CommandQueue::enqueueUnmapMemObject(cl_mem memobj,
void *mappedPtr,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event)
{
const Memory &memory = memobj->cast<Memory>();
const EventPtrs waitEvents = Event::Cast(numEventsInWaitList, eventWaitList);
rx::CLEventImpl::CreateFunc eventCreateFunc;
rx::CLEventImpl::CreateFunc *const eventCreateFuncPtr =
event != nullptr ? &eventCreateFunc : nullptr;
cl_int errorCode =
mImpl->enqueueUnmapMemObject(memory, mappedPtr, waitEvents, eventCreateFuncPtr);
CheckCreateEvent(*this, CL_COMMAND_UNMAP_MEM_OBJECT, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueMigrateMemObjects(cl_uint numMemObjects,
const cl_mem *memObjects,
MemMigrationFlags flags,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event)
{
MemoryPtrs memories;
memories.reserve(numMemObjects);
while (numMemObjects-- != 0u)
{
ASSERT(eventCreateFunc);
*event = Object::Create<Event>(errorCode, *this, CL_COMMAND_MAP_IMAGE, eventCreateFunc);
memories.emplace_back(&(*memObjects++)->cast<Memory>());
}
return map;
const EventPtrs waitEvents = Event::Cast(numEventsInWaitList, eventWaitList);
rx::CLEventImpl::CreateFunc eventCreateFunc;
rx::CLEventImpl::CreateFunc *const eventCreateFuncPtr =
event != nullptr ? &eventCreateFunc : nullptr;
cl_int errorCode =
mImpl->enqueueMigrateMemObjects(memories, flags, waitEvents, eventCreateFuncPtr);
CheckCreateEvent(*this, CL_COMMAND_MIGRATE_MEM_OBJECTS, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueNDRangeKernel(cl_kernel kernel,
cl_uint workDim,
const size_t *globalWorkOffset,
const size_t *globalWorkSize,
const size_t *localWorkSize,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event)
{
const Kernel &krnl = kernel->cast<Kernel>();
const EventPtrs waitEvents = Event::Cast(numEventsInWaitList, eventWaitList);
rx::CLEventImpl::CreateFunc eventCreateFunc;
rx::CLEventImpl::CreateFunc *const eventCreateFuncPtr =
event != nullptr ? &eventCreateFunc : nullptr;
cl_int errorCode = mImpl->enqueueNDRangeKernel(krnl, workDim, globalWorkOffset, globalWorkSize,
localWorkSize, waitEvents, eventCreateFuncPtr);
CheckCreateEvent(*this, CL_COMMAND_NDRANGE_KERNEL, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueTask(cl_kernel kernel,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event)
{
const Kernel &krnl = kernel->cast<Kernel>();
const EventPtrs waitEvents = Event::Cast(numEventsInWaitList, eventWaitList);
rx::CLEventImpl::CreateFunc eventCreateFunc;
rx::CLEventImpl::CreateFunc *const eventCreateFuncPtr =
event != nullptr ? &eventCreateFunc : nullptr;
cl_int errorCode = mImpl->enqueueTask(krnl, waitEvents, eventCreateFuncPtr);
CheckCreateEvent(*this, CL_COMMAND_TASK, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueNativeKernel(UserFunc userFunc,
void *args,
size_t cbArgs,
cl_uint numMemObjects,
const cl_mem *memList,
const void **argsMemLoc,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event)
{
std::vector<unsigned char> funcArgs;
BufferPtrs buffers;
std::vector<size_t> offsets;
if (numMemObjects != 0u)
{
// If argument memory block contains memory objects, make a copy.
funcArgs.resize(cbArgs);
std::memcpy(funcArgs.data(), args, cbArgs);
buffers.reserve(numMemObjects);
offsets.reserve(numMemObjects);
while (numMemObjects-- != 0u)
{
buffers.emplace_back(&(*memList++)->cast<Buffer>());
// Calc memory offset of cl_mem object in args.
offsets.emplace_back(static_cast<const char *>(*argsMemLoc++) -
static_cast<const char *>(args));
// Fetch location of cl_mem object in copied function argument memory block.
void *loc = &funcArgs[offsets.back()];
// Cast cl_mem object to cl::Buffer pointer in place.
*reinterpret_cast<Buffer **>(loc) = &(*reinterpret_cast<cl_mem *>(loc))->cast<Buffer>();
}
// Use copied argument memory block.
args = funcArgs.data();
}
const EventPtrs waitEvents = Event::Cast(numEventsInWaitList, eventWaitList);
rx::CLEventImpl::CreateFunc eventCreateFunc;
rx::CLEventImpl::CreateFunc *const eventCreateFuncPtr =
event != nullptr ? &eventCreateFunc : nullptr;
cl_int errorCode = mImpl->enqueueNativeKernel(userFunc, args, cbArgs, buffers, offsets,
waitEvents, eventCreateFuncPtr);
CheckCreateEvent(*this, CL_COMMAND_NATIVE_KERNEL, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueMarkerWithWaitList(cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event)
{
const EventPtrs waitEvents = Event::Cast(numEventsInWaitList, eventWaitList);
rx::CLEventImpl::CreateFunc eventCreateFunc;
rx::CLEventImpl::CreateFunc *const eventCreateFuncPtr =
event != nullptr ? &eventCreateFunc : nullptr;
cl_int errorCode = mImpl->enqueueMarkerWithWaitList(waitEvents, eventCreateFuncPtr);
CheckCreateEvent(*this, CL_COMMAND_MARKER, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueMarker(cl_event *event)
{
rx::CLEventImpl::CreateFunc eventCreateFunc;
cl_int errorCode = mImpl->enqueueMarker(eventCreateFunc);
CheckCreateEvent(*this, CL_COMMAND_MARKER, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueWaitForEvents(cl_uint numEvents, const cl_event *eventList)
{
return mImpl->enqueueWaitForEvents(Event::Cast(numEvents, eventList));
}
cl_int CommandQueue::enqueueBarrierWithWaitList(cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event)
{
const EventPtrs waitEvents = Event::Cast(numEventsInWaitList, eventWaitList);
rx::CLEventImpl::CreateFunc eventCreateFunc;
rx::CLEventImpl::CreateFunc *const eventCreateFuncPtr =
event != nullptr ? &eventCreateFunc : nullptr;
cl_int errorCode = mImpl->enqueueBarrierWithWaitList(waitEvents, eventCreateFuncPtr);
CheckCreateEvent(*this, CL_COMMAND_BARRIER, eventCreateFunc, event, errorCode);
return errorCode;
}
cl_int CommandQueue::enqueueBarrier()
{
return mImpl->enqueueBarrier();
}
cl_int CommandQueue::flush()
{
return mImpl->flush();
}
cl_int CommandQueue::finish()
{
return mImpl->finish();
}
CommandQueue::~CommandQueue()
......@@ -551,6 +701,12 @@ CommandQueue::~CommandQueue()
}
}
size_t CommandQueue::getDeviceIndex() const
{
return std::find(mContext->getDevices().cbegin(), mContext->getDevices().cend(), mDevice) -
mContext->getDevices().cbegin();
}
CommandQueue::CommandQueue(Context &context,
Device &device,
PropArray &&propArray,
......
src/libANGLE/CLCommandQueue.h
......@@ -187,6 +187,60 @@ class CommandQueue final : public _cl_command_queue, public Object
cl_event *event,
cl_int &errorCode);
cl_int enqueueUnmapMemObject(cl_mem memobj,
void *mappedPtr,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event);
cl_int enqueueMigrateMemObjects(cl_uint numMemObjects,
const cl_mem *memObjects,
MemMigrationFlags flags,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event);
cl_int enqueueNDRangeKernel(cl_kernel kernel,
cl_uint workDim,
const size_t *globalWorkOffset,
const size_t *globalWorkSize,
const size_t *localWorkSize,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event);
cl_int enqueueTask(cl_kernel kernel,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event);
cl_int enqueueNativeKernel(UserFunc userFunc,
void *args,
size_t cbArgs,
cl_uint numMemObjects,
const cl_mem *memList,
const void **argsMemLoc,
cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event);
cl_int enqueueMarkerWithWaitList(cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event);
cl_int enqueueMarker(cl_event *event);
cl_int enqueueWaitForEvents(cl_uint numEvents, const cl_event *eventList);
cl_int enqueueBarrierWithWaitList(cl_uint numEventsInWaitList,
const cl_event *eventWaitList,
cl_event *event);
cl_int enqueueBarrier();
cl_int flush();
cl_int finish();
public:
using PropArray = std::vector<cl_queue_properties>;
......@@ -198,6 +252,9 @@ class CommandQueue final : public _cl_command_queue, public Object
const Context &getContext() const;
const Device &getDevice() const;
// Get index of device in the context.
size_t getDeviceIndex() const;
CommandQueueProperties getProperties() const;
bool isOnHost() const;
bool isOnDevice() const;
......
src/libANGLE/CLContext.h
......@@ -172,9 +172,7 @@ inline const DevicePtrs &Context::getDevices() const
inline bool Context::hasDevice(const _cl_device_id *device) const
{
return std::find_if(mDevices.cbegin(), mDevices.cend(), [=](const DevicePtr &ptr) {
return ptr.get() == device;
}) != mDevices.cend();
return std::find(mDevices.cbegin(), mDevices.cend(), device) != mDevices.cend();
}
template <typename T>
......
src/libANGLE/CLDevice.cpp
......@@ -108,7 +108,6 @@ cl_int Device::getInfo(DeviceInfo name, size_t valueSize, void *value, size_t *v
case DeviceInfo::GlobalMemSize:
case DeviceInfo::MaxConstantBufferSize:
case DeviceInfo::LocalMemSize:
case DeviceInfo::ExecutionCapabilities:
case DeviceInfo::QueueOnHostProperties:
case DeviceInfo::QueueOnDeviceProperties:
case DeviceInfo::PartitionAffinityDomain:
......@@ -224,6 +223,10 @@ cl_int Device::getInfo(DeviceInfo name, size_t valueSize, void *value, size_t *v
copyValue = &mInfo.mMemBaseAddrAlign;
copySize = sizeof(mInfo.mMemBaseAddrAlign);
break;
case DeviceInfo::ExecutionCapabilities:
copyValue = &mInfo.mExecCapabilities;
copySize = sizeof(mInfo.mExecCapabilities);
break;
case DeviceInfo::QueueOnDeviceMaxSize:
copyValue = &mInfo.mQueueOnDeviceMaxSize;
copySize = sizeof(mInfo.mQueueOnDeviceMaxSize);
......
src/libANGLE/CLKernel.cpp
......@@ -85,7 +85,7 @@ cl_int Kernel::getWorkGroupInfo(cl_device_id device,
if (device != nullptr)
{
const DevicePtrs &devices = mProgram->getContext().getDevices();
while (index < devices.size() && devices[index].get() != device)
while (index < devices.size() && devices[index] != device)
{
++index;
}
......
src/libANGLE/CLProgram.h
......@@ -118,9 +118,7 @@ inline const DevicePtrs &Program::getDevices() const
inline bool Program::hasDevice(const _cl_device_id *device) const
{
return std::find_if(mDevices.cbegin(), mDevices.cend(), [=](const DevicePtr &ptr) {
return ptr.get() == device;
}) != mDevices.cend();
return std::find(mDevices.cbegin(), mDevices.cend(), device) != mDevices.cend();
}
inline bool Program::isBuilding() const
......
src/libANGLE/CLRefPointer.h
......@@ -127,6 +127,42 @@ bool operator!=(nullptr_t, const RefPointer<T> &ptr) noexcept
return ptr.get() != nullptr;
}
template <typename T, typename U>
bool operator==(const RefPointer<T> &left, const RefPointer<U> &right) noexcept
{
return left.get() == right.get();
}
template <typename T, typename U>
bool operator!=(const RefPointer<T> &left, const RefPointer<U> &right) noexcept
{
return left.get() != right.get();
}
template <typename T, typename U>
bool operator==(const RefPointer<T> &left, const U *right) noexcept
{
return left.get() == right;
}
template <typename T, typename U>
bool operator==(const T *left, const RefPointer<U> &right) noexcept
{
return left == right.get();
}
template <typename T, typename U>
bool operator!=(const RefPointer<T> &left, const U *right) noexcept
{
return left.get() != right;
}
template <typename T, typename U>
bool operator!=(const T *left, const RefPointer<U> &right) noexcept
{
return left != right.get();
}
} // namespace cl
#endif // LIBANGLE_CLREFPOINTER_H_
src/libANGLE/CLtypes.h
......@@ -40,6 +40,7 @@ class Platform;
class Program;
class Sampler;
using BufferPtr = RefPointer<Buffer>;
using CommandQueuePtr = RefPointer<CommandQueue>;
using ContextPtr = RefPointer<Context>;
using DevicePtr = RefPointer<Device>;
......@@ -50,9 +51,11 @@ using PlatformPtr = RefPointer<Platform>;
using ProgramPtr = RefPointer<Program>;
using SamplerPtr = RefPointer<Sampler>;
using BufferPtrs = std::vector<BufferPtr>;
using DevicePtrs = std::vector<DevicePtr>;
using EventPtrs = std::vector<EventPtr>;
using KernelPtrs = std::vector<KernelPtr>;
using MemoryPtrs = std::vector<MemoryPtr>;
using PlatformPtrs = std::vector<PlatformPtr>;
using ProgramPtrs = std::vector<ProgramPtr>;
......
src/libANGLE/renderer/CLCommandQueueImpl.h
......@@ -164,6 +164,51 @@ class CLCommandQueueImpl : angle::NonCopyable
CLEventImpl::CreateFunc *eventCreateFunc,
cl_int &errorCode) = 0;
virtual cl_int enqueueUnmapMemObject(const cl::Memory &memory,
void *mappedPtr,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) = 0;
virtual cl_int enqueueMigrateMemObjects(const cl::MemoryPtrs &memObjects,
cl::MemMigrationFlags flags,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) = 0;
virtual cl_int enqueueNDRangeKernel(const cl::Kernel &kernel,
cl_uint workDim,
const size_t *globalWorkOffset,
const size_t *globalWorkSize,
const size_t *localWorkSize,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) = 0;
virtual cl_int enqueueTask(const cl::Kernel &kernel,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) = 0;
virtual cl_int enqueueNativeKernel(cl::UserFunc userFunc,
void *args,
size_t cbArgs,
const cl::BufferPtrs &buffers,
const std::vector<size_t> bufferPtrOffsets,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) = 0;
virtual cl_int enqueueMarkerWithWaitList(const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) = 0;
virtual cl_int enqueueMarker(CLEventImpl::CreateFunc &eventCreateFunc) = 0;
virtual cl_int enqueueWaitForEvents(const cl::EventPtrs &events) = 0;
virtual cl_int enqueueBarrierWithWaitList(const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) = 0;
virtual cl_int enqueueBarrier() = 0;
virtual cl_int flush() = 0;
virtual cl_int finish() = 0;
protected:
const cl::CommandQueue &mCommandQueue;
};
......
src/libANGLE/renderer/CLDeviceImpl.h
......@@ -53,7 +53,8 @@ class CLDeviceImpl : angle::NonCopyable
cl_uint mImagePitchAlignment = 0u;
cl_uint mImageBaseAddressAlignment = 0u;
cl_uint mMemBaseAddrAlign = 0u;
cl_uint mQueueOnDeviceMaxSize = 0u;
cl::DeviceExecCapabilities mExecCapabilities;
cl_uint mQueueOnDeviceMaxSize = 0u;
std::string mBuiltInKernels;
NameVersionVector mBuiltInKernelsWithVersion;
std::string mVersionStr;
......
src/libANGLE/renderer/cl/CLCommandQueueCL.cpp
......@@ -8,14 +8,32 @@
#include "libANGLE/renderer/cl/CLCommandQueueCL.h"
#include "libANGLE/renderer/cl/CLEventCL.h"
#include "libANGLE/renderer/cl/CLKernelCL.h"
#include "libANGLE/renderer/cl/CLMemoryCL.h"
#include "libANGLE/CLBuffer.h"
#include "libANGLE/CLImage.h"
#include "libANGLE/CLKernel.h"
#include "libANGLE/CLMemory.h"
namespace rx
{
namespace
{
void CheckCreateEvent(cl_int errorCode, cl_event nativeEvent, CLEventImpl::CreateFunc *createFunc)
{
if (errorCode == CL_SUCCESS && createFunc != nullptr)
{
*createFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
}
} // namespace
CLCommandQueueCL::CLCommandQueueCL(const cl::CommandQueue &commandQueue, cl_command_queue native)
: CLCommandQueueImpl(commandQueue), mNative(native)
{}
......@@ -54,12 +72,7 @@ cl_int CLCommandQueueCL::enqueueReadBuffer(const cl::Buffer &buffer,
mNative->getDispatch().clEnqueueReadBuffer(mNative, nativeBuffer, block, offset, size, ptr,
numEvents, nativeEventsPtr, nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -83,12 +96,7 @@ cl_int CLCommandQueueCL::enqueueWriteBuffer(const cl::Buffer &buffer,
mNative->getDispatch().clEnqueueWriteBuffer(mNative, nativeBuffer, block, offset, size, ptr,
numEvents, nativeEventsPtr, nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -118,12 +126,7 @@ cl_int CLCommandQueueCL::enqueueReadBufferRect(const cl::Buffer &buffer,
bufferSlicePitch, hostRowPitch, hostSlicePitch, ptr, numEvents, nativeEventsPtr,
nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -153,12 +156,7 @@ cl_int CLCommandQueueCL::enqueueWriteBufferRect(const cl::Buffer &buffer,
bufferSlicePitch, hostRowPitch, hostSlicePitch, ptr, numEvents, nativeEventsPtr,
nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -182,12 +180,7 @@ cl_int CLCommandQueueCL::enqueueCopyBuffer(const cl::Buffer &srcBuffer,
mNative, nativeSrc, nativeDst, srcOffset, dstOffset, size, numEvents, nativeEventsPtr,
nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -215,12 +208,7 @@ cl_int CLCommandQueueCL::enqueueCopyBufferRect(const cl::Buffer &srcBuffer,
mNative, nativeSrc, nativeDst, srcOrigin, dstOrigin, region, srcRowPitch, srcSlicePitch,
dstRowPitch, dstSlicePitch, numEvents, nativeEventsPtr, nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -243,12 +231,7 @@ cl_int CLCommandQueueCL::enqueueFillBuffer(const cl::Buffer &buffer,
mNative, nativeBuffer, pattern, patternSize, offset, size, numEvents, nativeEventsPtr,
nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -273,12 +256,7 @@ void *CLCommandQueueCL::enqueueMapBuffer(const cl::Buffer &buffer,
mNative, nativeBuffer, block, mapFlags.get(), offset, size, numEvents, nativeEventsPtr,
nativeEventPtr, &errorCode);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return map;
}
......@@ -304,12 +282,7 @@ cl_int CLCommandQueueCL::enqueueReadImage(const cl::Image &image,
mNative, nativeImage, block, origin, region, rowPitch, slicePitch, ptr, numEvents,
nativeEventsPtr, nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -335,12 +308,7 @@ cl_int CLCommandQueueCL::enqueueWriteImage(const cl::Image &image,
mNative, nativeImage, block, origin, region, inputRowPitch, inputSlicePitch, ptr, numEvents,
nativeEventsPtr, nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -364,12 +332,7 @@ cl_int CLCommandQueueCL::enqueueCopyImage(const cl::Image &srcImage,
mNative, nativeSrc, nativeDst, srcOrigin, dstOrigin, region, numEvents, nativeEventsPtr,
nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -391,12 +354,7 @@ cl_int CLCommandQueueCL::enqueueFillImage(const cl::Image &image,
mNative->getDispatch().clEnqueueFillImage(mNative, nativeImage, fillColor, origin, region,
numEvents, nativeEventsPtr, nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -420,12 +378,7 @@ cl_int CLCommandQueueCL::enqueueCopyImageToBuffer(const cl::Image &srcImage,
mNative, nativeSrc, nativeDst, srcOrigin, region, dstOffset, numEvents, nativeEventsPtr,
nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -449,12 +402,7 @@ cl_int CLCommandQueueCL::enqueueCopyBufferToImage(const cl::Buffer &srcBuffer,
mNative, nativeSrc, nativeDst, srcOffset, dstOrigin, region, numEvents, nativeEventsPtr,
nativeEventPtr);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
......@@ -481,13 +429,220 @@ void *CLCommandQueueCL::enqueueMapImage(const cl::Image &image,
mNative, nativeImage, block, mapFlags.get(), origin, region, imageRowPitch, imageSlicePitch,
numEvents, nativeEventsPtr, nativeEventPtr, &errorCode);
if (errorCode == CL_SUCCESS && eventCreateFunc != nullptr)
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return map;
}
cl_int CLCommandQueueCL::enqueueUnmapMemObject(const cl::Memory &memory,
void *mappedPtr,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc)
{
const cl_mem nativeMemory = memory.getImpl<CLMemoryCL>().getNative();
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(waitEvents);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
const cl_event *const nativeEventsPtr = nativeEvents.empty() ? nullptr : nativeEvents.data();
cl_event nativeEvent = nullptr;
cl_event *const nativeEventPtr = eventCreateFunc != nullptr ? &nativeEvent : nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueUnmapMemObject(
mNative, nativeMemory, mappedPtr, numEvents, nativeEventsPtr, nativeEventPtr);
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
cl_int CLCommandQueueCL::enqueueMigrateMemObjects(const cl::MemoryPtrs &memObjects,
cl::MemMigrationFlags flags,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc)
{
std::vector<cl_mem> nativeMemories;
nativeMemories.reserve(memObjects.size());
for (const cl::MemoryPtr &memory : memObjects)
{
nativeMemories.emplace_back(memory->getImpl<CLMemoryCL>().getNative());
}
const cl_uint numMemories = static_cast<cl_uint>(nativeMemories.size());
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(waitEvents);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
const cl_event *const nativeEventsPtr = nativeEvents.empty() ? nullptr : nativeEvents.data();
cl_event nativeEvent = nullptr;
cl_event *const nativeEventPtr = eventCreateFunc != nullptr ? &nativeEvent : nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueMigrateMemObjects(
mNative, numMemories, nativeMemories.data(), flags.get(), numEvents, nativeEventsPtr,
nativeEventPtr);
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
cl_int CLCommandQueueCL::enqueueNDRangeKernel(const cl::Kernel &kernel,
cl_uint workDim,
const size_t *globalWorkOffset,
const size_t *globalWorkSize,
const size_t *localWorkSize,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc)
{
const cl_kernel nativeKernel = kernel.getImpl<CLKernelCL>().getNative();
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(waitEvents);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
const cl_event *const nativeEventsPtr = nativeEvents.empty() ? nullptr : nativeEvents.data();
cl_event nativeEvent = nullptr;
cl_event *const nativeEventPtr = eventCreateFunc != nullptr ? &nativeEvent : nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueNDRangeKernel(
mNative, nativeKernel, workDim, globalWorkOffset, globalWorkSize, localWorkSize, numEvents,
nativeEventsPtr, nativeEventPtr);
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
cl_int CLCommandQueueCL::enqueueTask(const cl::Kernel &kernel,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc)
{
const cl_kernel nativeKernel = kernel.getImpl<CLKernelCL>().getNative();
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(waitEvents);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
const cl_event *const nativeEventsPtr = nativeEvents.empty() ? nullptr : nativeEvents.data();
cl_event nativeEvent = nullptr;
cl_event *const nativeEventPtr = eventCreateFunc != nullptr ? &nativeEvent : nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueTask(mNative, nativeKernel, numEvents,
nativeEventsPtr, nativeEventPtr);
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
cl_int CLCommandQueueCL::enqueueNativeKernel(cl::UserFunc userFunc,
void *args,
size_t cbArgs,
const cl::BufferPtrs &buffers,
const std::vector<size_t> bufferPtrOffsets,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc)
{
std::vector<unsigned char> funcArgs;
std::vector<const void *> locs;
if (!bufferPtrOffsets.empty())
{
// If argument memory block contains buffers, make a copy.
funcArgs.resize(cbArgs);
std::memcpy(funcArgs.data(), args, cbArgs);
locs.reserve(bufferPtrOffsets.size());
for (size_t offset : bufferPtrOffsets)
{
// Fetch location of buffer in copied function argument memory block.
void *const loc = &funcArgs[offset];
locs.emplace_back(loc);
// Cast cl::Buffer to native cl_mem object in place.
cl::Buffer *const buffer = *reinterpret_cast<cl::Buffer **>(loc);
*reinterpret_cast<cl_mem *>(loc) = buffer->getImpl<CLMemoryCL>().getNative();
}
// Use copied argument memory block.
args = funcArgs.data();
}
std::vector<cl_mem> nativeBuffers;
nativeBuffers.reserve(buffers.size());
for (const cl::BufferPtr &buffer : buffers)
{
nativeBuffers.emplace_back(buffer->getImpl<CLMemoryCL>().getNative());
}
const cl_uint numBuffers = static_cast<cl_uint>(nativeBuffers.size());
const cl_mem *const nativeBuffersPtr = nativeBuffers.empty() ? nullptr : nativeBuffers.data();
const void **const locsPtr = locs.empty() ? nullptr : locs.data();
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(waitEvents);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
const cl_event *const nativeEventsPtr = nativeEvents.empty() ? nullptr : nativeEvents.data();
cl_event nativeEvent = nullptr;
cl_event *const nativeEventPtr = eventCreateFunc != nullptr ? &nativeEvent : nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueNativeKernel(
mNative, userFunc, args, cbArgs, numBuffers, nativeBuffersPtr, locsPtr, numEvents,
nativeEventsPtr, nativeEventPtr);
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
cl_int CLCommandQueueCL::enqueueMarkerWithWaitList(const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc)
{
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(waitEvents);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
const cl_event *const nativeEventsPtr = nativeEvents.empty() ? nullptr : nativeEvents.data();
cl_event nativeEvent = nullptr;
cl_event *const nativeEventPtr = eventCreateFunc != nullptr ? &nativeEvent : nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueMarkerWithWaitList(
mNative, numEvents, nativeEventsPtr, nativeEventPtr);
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
cl_int CLCommandQueueCL::enqueueMarker(CLEventImpl::CreateFunc &eventCreateFunc)
{
cl_event nativeEvent = nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueMarker(mNative, &nativeEvent);
if (errorCode == CL_SUCCESS)
{
*eventCreateFunc = [nativeEvent](const cl::Event &event) {
eventCreateFunc = [nativeEvent](const cl::Event &event) {
return CLEventImpl::Ptr(new CLEventCL(event, nativeEvent));
};
}
return map;
return errorCode;
}
cl_int CLCommandQueueCL::enqueueWaitForEvents(const cl::EventPtrs &events)
{
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(events);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
return mNative->getDispatch().clEnqueueWaitForEvents(mNative, numEvents, nativeEvents.data());
}
cl_int CLCommandQueueCL::enqueueBarrierWithWaitList(const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc)
{
const std::vector<cl_event> nativeEvents = CLEventCL::Cast(waitEvents);
const cl_uint numEvents = static_cast<cl_uint>(nativeEvents.size());
const cl_event *const nativeEventsPtr = nativeEvents.empty() ? nullptr : nativeEvents.data();
cl_event nativeEvent = nullptr;
cl_event *const nativeEventPtr = eventCreateFunc != nullptr ? &nativeEvent : nullptr;
const cl_int errorCode = mNative->getDispatch().clEnqueueBarrierWithWaitList(
mNative, numEvents, nativeEventsPtr, nativeEventPtr);
CheckCreateEvent(errorCode, nativeEvent, eventCreateFunc);
return errorCode;
}
cl_int CLCommandQueueCL::enqueueBarrier()
{
return mNative->getDispatch().clEnqueueBarrier(mNative);
}
cl_int CLCommandQueueCL::flush()
{
return mNative->getDispatch().clFlush(mNative);
}
cl_int CLCommandQueueCL::finish()
{
return mNative->getDispatch().clFinish(mNative);
}
} // namespace rx
src/libANGLE/renderer/cl/CLCommandQueueCL.h
......@@ -163,6 +163,51 @@ class CLCommandQueueCL : public CLCommandQueueImpl
CLEventImpl::CreateFunc *eventCreateFunc,
cl_int &errorCode) override;
cl_int enqueueUnmapMemObject(const cl::Memory &memory,
void *mappedPtr,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) override;
cl_int enqueueMigrateMemObjects(const cl::MemoryPtrs &memObjects,
cl::MemMigrationFlags flags,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) override;
cl_int enqueueNDRangeKernel(const cl::Kernel &kernel,
cl_uint workDim,
const size_t *globalWorkOffset,
const size_t *globalWorkSize,
const size_t *localWorkSize,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) override;
cl_int enqueueTask(const cl::Kernel &kernel,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) override;
cl_int enqueueNativeKernel(cl::UserFunc userFunc,
void *args,
size_t cbArgs,
const cl::BufferPtrs &buffers,
const std::vector<size_t> bufferPtrOffsets,
const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) override;
cl_int enqueueMarkerWithWaitList(const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) override;
cl_int enqueueMarker(CLEventImpl::CreateFunc &eventCreateFunc) override;
cl_int enqueueWaitForEvents(const cl::EventPtrs &events) override;
cl_int enqueueBarrierWithWaitList(const cl::EventPtrs &waitEvents,
CLEventImpl::CreateFunc *eventCreateFunc) override;
cl_int enqueueBarrier() override;
cl_int flush() override;
cl_int finish() override;
private:
const cl_command_queue mNative;
};
......
src/libANGLE/renderer/cl/CLDeviceCL.cpp
......@@ -90,7 +90,8 @@ CLDeviceImpl::Info CLDeviceCL::createInfo(cl::DeviceType type) const
!GetDeviceInfo(mNative, cl::DeviceInfo::Image3D_MaxWidth, info.mImage3D_MaxWidth) ||
!GetDeviceInfo(mNative, cl::DeviceInfo::Image3D_MaxHeight, info.mImage3D_MaxHeight) ||
!GetDeviceInfo(mNative, cl::DeviceInfo::Image3D_MaxDepth, info.mImage3D_MaxDepth) ||
!GetDeviceInfo(mNative, cl::DeviceInfo::MemBaseAddrAlign, info.mMemBaseAddrAlign))
!GetDeviceInfo(mNative, cl::DeviceInfo::MemBaseAddrAlign, info.mMemBaseAddrAlign) ||
!GetDeviceInfo(mNative, cl::DeviceInfo::ExecutionCapabilities, info.mExecCapabilities))
{
return Info{};
}
......
src/libANGLE/validationCL.cpp
......@@ -1507,11 +1507,21 @@ cl_int ValidateGetEventProfilingInfo(cl_event event,
cl_int ValidateFlush(cl_command_queue command_queue)
{
// CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
if (!CommandQueue::IsValid(command_queue) || !command_queue->cast<CommandQueue>().isOnHost())
{
return CL_INVALID_COMMAND_QUEUE;
}
return CL_SUCCESS;
}
cl_int ValidateFinish(cl_command_queue command_queue)
{
// CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
if (!CommandQueue::IsValid(command_queue) || !command_queue->cast<CommandQueue>().isOnHost())
{
return CL_INVALID_COMMAND_QUEUE;
}
return CL_SUCCESS;
}
......@@ -1743,7 +1753,7 @@ cl_int ValidateEnqueueCopyImageToBuffer(cl_command_queue command_queue,
const Buffer &dst = dst_buffer->cast<Buffer>();
// CL_INVALID_MEM_OBJECT if src_image is a 1D image buffer object created from dst_buffer.
if (src.getType() == CL_MEM_OBJECT_IMAGE1D_BUFFER && src.getParent().get() == &dst)
if (src.getType() == CL_MEM_OBJECT_IMAGE1D_BUFFER && src.getParent() == &dst)
{
return CL_INVALID_MEM_OBJECT;
}
......@@ -1791,7 +1801,7 @@ cl_int ValidateEnqueueCopyBufferToImage(cl_command_queue command_queue,
const Image &dst = dst_image->cast<Image>();
// CL_INVALID_MEM_OBJECT if dst_image is a 1D image buffer object created from src_buffer.
if (dst.getType() == CL_MEM_OBJECT_IMAGE1D_BUFFER && dst.getParent().get() == &src)
if (dst.getType() == CL_MEM_OBJECT_IMAGE1D_BUFFER && dst.getParent() == &src)
{
return CL_INVALID_MEM_OBJECT;
}
......@@ -1908,6 +1918,27 @@ cl_int ValidateEnqueueUnmapMemObject(cl_command_queue command_queue,
const cl_event *event_wait_list,
const cl_event *event)
{
ANGLE_TRY(ValidateCommandQueueAndEventWaitList(command_queue, false, num_events_in_wait_list,
event_wait_list));
const CommandQueue &queue = command_queue->cast<CommandQueue>();
// CL_INVALID_MEM_OBJECT if memobj is not a valid memory object or is a pipe object.
if (!Memory::IsValid(memobj))
{
return CL_INVALID_MEM_OBJECT;
}
const Memory &memory = memobj->cast<Memory>();
if (memory.getType() == CL_MEM_OBJECT_PIPE)
{
return CL_INVALID_MEM_OBJECT;
}
// CL_INVALID_CONTEXT if context associated with command_queue and memobj are not the same.
if (&queue.getContext() != &memory.getContext())
{
return CL_INVALID_CONTEXT;
}
return CL_SUCCESS;
}
......@@ -1921,6 +1952,81 @@ cl_int ValidateEnqueueNDRangeKernel(cl_command_queue command_queue,
const cl_event *event_wait_list,
const cl_event *event)
{
ANGLE_TRY(ValidateCommandQueueAndEventWaitList(command_queue, false, num_events_in_wait_list,
event_wait_list));
const CommandQueue &queue = command_queue->cast<CommandQueue>();
const Device &device = queue.getDevice();
// CL_INVALID_KERNEL if kernel is not a valid kernel object.
if (!Kernel::IsValid(kernel))
{
return CL_INVALID_KERNEL;
}
const Kernel &krnl = kernel->cast<Kernel>();
// CL_INVALID_CONTEXT if context associated with command_queue and kernel are not the same.
if (&queue.getContext() != &krnl.getProgram().getContext())
{
return CL_INVALID_CONTEXT;
}
// CL_INVALID_WORK_DIMENSION if work_dim is not a valid value.
if (work_dim == 0u || work_dim > device.getInfo().mMaxWorkItemSizes.size())
{
return CL_INVALID_WORK_DIMENSION;
}
// CL_INVALID_GLOBAL_OFFSET if global_work_offset is non-NULL before version 1.1.
if (!queue.getContext().getPlatform().isVersionOrNewer(1u, 1u) && global_work_offset != nullptr)
{
return CL_INVALID_GLOBAL_OFFSET;
}
// CL_INVALID_GLOBAL_WORK_SIZE if global_work_size is NULL or if any of the values
// specified in global_work_size[0] ... global_work_size[work_dim - 1] are 0.
// Returning this error code under these circumstances is deprecated by version 2.1.
if (!queue.getContext().getPlatform().isVersionOrNewer(2u, 1u))
{
if (global_work_size == nullptr)
{
return CL_INVALID_GLOBAL_WORK_SIZE;
}
for (cl_uint dim = 0u; dim < work_dim; ++dim)
{
if (global_work_size[dim] == 0u)
{
return CL_INVALID_GLOBAL_WORK_SIZE;
}
}
}
if (local_work_size != nullptr)
{
size_t numWorkItems = 1u; // Initialize with neutral element for multiplication
// CL_INVALID_WORK_ITEM_SIZE if the number of work-items specified
// in any of local_work_size[0] ... local_work_size[work_dim - 1]
// is greater than the corresponding values specified by
// CL_DEVICE_MAX_WORK_ITEM_SIZES[0] ... CL_DEVICE_MAX_WORK_ITEM_SIZES[work_dim - 1].
for (cl_uint dim = 0u; dim < work_dim; ++dim)
{
if (local_work_size[dim] > device.getInfo().mMaxWorkItemSizes[dim])
{
return CL_INVALID_WORK_ITEM_SIZE;
}
numWorkItems *= local_work_size[dim];
}
// CL_INVALID_WORK_GROUP_SIZE if local_work_size is specified
// and the total number of work-items in the work-group computed as
// local_work_size[0] x ... local_work_size[work_dim - 1] is greater than the value
// specified by CL_KERNEL_WORK_GROUP_SIZE in the Kernel Object Device Queries table.
if (numWorkItems > krnl.getInfo().mWorkGroups[queue.getDeviceIndex()].mWorkGroupSize)
{
return CL_INVALID_WORK_GROUP_SIZE;
}
}
return CL_SUCCESS;
}
......@@ -1935,6 +2041,67 @@ cl_int ValidateEnqueueNativeKernel(cl_command_queue command_queue,
const cl_event *event_wait_list,
const cl_event *event)
{
ANGLE_TRY(ValidateCommandQueueAndEventWaitList(command_queue, false, num_events_in_wait_list,
event_wait_list));
const CommandQueue &queue = command_queue->cast<CommandQueue>();
// CL_INVALID_OPERATION if the device associated with command_queue
// cannot execute the native kernel.
if (queue.getDevice().getInfo().mExecCapabilities.isNotSet(CL_EXEC_NATIVE_KERNEL))
{
return CL_INVALID_OPERATION;
}
// CL_INVALID_VALUE if user_func is NULL.
if (user_func == nullptr)
{
return CL_INVALID_VALUE;
}
if (args == nullptr)
{
// CL_INVALID_VALUE if args is a NULL value and cb_args > 0 or num_mem_objects > 0.
if (cb_args > 0u || num_mem_objects > 0u)
{
return CL_INVALID_VALUE;
}
}
else
{
// CL_INVALID_VALUE if args is not NULL and cb_args is 0.
if (cb_args == 0u)
{
return CL_INVALID_VALUE;
}
}
if (num_mem_objects == 0u)
{
// CL_INVALID_VALUE if num_mem_objects = 0 and mem_list or args_mem_loc are not NULL.
if (mem_list != nullptr || args_mem_loc != nullptr)
{
return CL_INVALID_VALUE;
}
}
else
{
// CL_INVALID_VALUE if num_mem_objects > 0 and mem_list or args_mem_loc are NULL.
if (mem_list == nullptr || args_mem_loc == nullptr)
{
return CL_INVALID_VALUE;
}
// CL_INVALID_MEM_OBJECT if one or more memory objects
// specified in mem_list are not valid or are not buffer objects.
while (num_mem_objects-- != 0u)
{
if (!Buffer::IsValid(*mem_list++))
{
return CL_INVALID_MEM_OBJECT;
}
}
}
return CL_SUCCESS;
}
......@@ -1990,6 +2157,18 @@ cl_int ValidateCreateImage3D(cl_context context,
cl_int ValidateEnqueueMarker(cl_command_queue command_queue, const cl_event *event)
{
// CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
if (!CommandQueue::IsValid(command_queue) || !command_queue->cast<CommandQueue>().isOnHost())
{
return CL_INVALID_COMMAND_QUEUE;
}
// CL_INVALID_VALUE if event is NULL.
if (event == nullptr)
{
return CL_INVALID_VALUE;
}
return CL_SUCCESS;
}
......@@ -1997,11 +2176,49 @@ cl_int ValidateEnqueueWaitForEvents(cl_command_queue command_queue,
cl_uint num_events,
const cl_event *event_list)
{
// CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
if (!CommandQueue::IsValid(command_queue))
{
return CL_INVALID_COMMAND_QUEUE;
}
const CommandQueue &queue = command_queue->cast<CommandQueue>();
if (!queue.isOnHost())
{
return CL_INVALID_COMMAND_QUEUE;
}
// CL_INVALID_VALUE if num_events is 0 or event_list is NULL.
if (num_events == 0u || event_list == nullptr)
{
return CL_INVALID_VALUE;
}
while (num_events-- != 0u)
{
// The documentation for invalid events is missing.
if (!Event::IsValid(*event_list))
{
return CL_INVALID_VALUE;
}
// CL_INVALID_CONTEXT if context associated with command_queue
// and events in event_list are not the same.
if (&queue.getContext() != &(*event_list++)->cast<Event>().getContext())
{
return CL_INVALID_CONTEXT;
}
}
return CL_SUCCESS;
}
cl_int ValidateEnqueueBarrier(cl_command_queue command_queue)
{
// CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
if (!CommandQueue::IsValid(command_queue) || !command_queue->cast<CommandQueue>().isOnHost())
{
return CL_INVALID_COMMAND_QUEUE;
}
return CL_SUCCESS;
}
......@@ -2075,6 +2292,22 @@ cl_int ValidateEnqueueTask(cl_command_queue command_queue,
const cl_event *event_wait_list,
const cl_event *event)
{
ANGLE_TRY(ValidateCommandQueueAndEventWaitList(command_queue, false, num_events_in_wait_list,
event_wait_list));
// CL_INVALID_KERNEL if kernel is not a valid kernel object.
if (!Kernel::IsValid(kernel))
{
return CL_INVALID_KERNEL;
}
// CL_INVALID_CONTEXT if context associated with command_queue and kernel are not the same.
if (&command_queue->cast<CommandQueue>().getContext() !=
&kernel->cast<Kernel>().getProgram().getContext())
{
return CL_INVALID_CONTEXT;
}
return CL_SUCCESS;
}
......@@ -2856,6 +3089,45 @@ cl_int ValidateEnqueueMigrateMemObjects(cl_command_queue command_queue,
const cl_event *event_wait_list,
const cl_event *event)
{
ANGLE_TRY(ValidateCommandQueueAndEventWaitList(command_queue, false, num_events_in_wait_list,
event_wait_list));
const CommandQueue &queue = command_queue->cast<CommandQueue>();
if (!queue.getContext().getPlatform().isVersionOrNewer(1u, 2u))
{
return CL_INVALID_COMMAND_QUEUE;
}
// CL_INVALID_VALUE if num_mem_objects is zero or if mem_objects is NULL.
if (num_mem_objects == 0u || mem_objects == nullptr)
{
return CL_INVALID_VALUE;
}
while (num_mem_objects-- != 0u)
{
// CL_INVALID_MEM_OBJECT if any of the memory objects
// in mem_objects is not a valid memory object.
if (!Memory::IsValid(*mem_objects))
{
return CL_INVALID_MEM_OBJECT;
}
// CL_INVALID_CONTEXT if the context associated with command_queue
// and memory objects in mem_objects are not the same.
if (&queue.getContext() != &(*mem_objects++)->cast<Memory>().getContext())
{
return CL_INVALID_CONTEXT;
}
}
// CL_INVALID_VALUE if flags is not 0 or is not any of the values described in the table.
const MemMigrationFlags allowedFlags(CL_MIGRATE_MEM_OBJECT_HOST |
CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED);
if (flags.hasOtherBitsThan(allowedFlags))
{
return CL_INVALID_VALUE;
}
return CL_SUCCESS;
}
......@@ -2864,6 +3136,12 @@ cl_int ValidateEnqueueMarkerWithWaitList(cl_command_queue command_queue,
const cl_event *event_wait_list,
const cl_event *event)
{
ANGLE_TRY(ValidateCommandQueueAndEventWaitList(command_queue, false, num_events_in_wait_list,
event_wait_list));
if (!command_queue->cast<CommandQueue>().getContext().getPlatform().isVersionOrNewer(1u, 2u))
{
return CL_INVALID_COMMAND_QUEUE;
}
return CL_SUCCESS;
}
......@@ -2872,6 +3150,12 @@ cl_int ValidateEnqueueBarrierWithWaitList(cl_command_queue command_queue,
const cl_event *event_wait_list,
const cl_event *event)
{
ANGLE_TRY(ValidateCommandQueueAndEventWaitList(command_queue, false, num_events_in_wait_list,
event_wait_list));
if (!command_queue->cast<CommandQueue>().getContext().getPlatform().isVersionOrNewer(1u, 2u))
{
return CL_INVALID_COMMAND_QUEUE;
}
return CL_SUCCESS;
}
......
src/libGLESv2/cl_stubs.cpp
......@@ -694,14 +694,12 @@ cl_int GetEventProfilingInfo(cl_event event,
cl_int Flush(cl_command_queue command_queue)
{
WARN_NOT_SUPPORTED(Flush);
return 0;
return command_queue->cast<CommandQueue>().flush();
}
cl_int Finish(cl_command_queue command_queue)
{
WARN_NOT_SUPPORTED(Finish);
return 0;
return command_queue->cast<CommandQueue>().finish();
}
cl_int EnqueueReadBuffer(cl_command_queue command_queue,
......@@ -956,8 +954,8 @@ cl_int EnqueueUnmapMemObject(cl_command_queue command_queue,
const cl_event *event_wait_list,
cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueUnmapMemObject);
return 0;
return command_queue->cast<CommandQueue>().enqueueUnmapMemObject(
memobj, mapped_ptr, num_events_in_wait_list, event_wait_list, event);
}
cl_int EnqueueMigrateMemObjects(cl_command_queue command_queue,
......@@ -968,8 +966,8 @@ cl_int EnqueueMigrateMemObjects(cl_command_queue command_queue,
const cl_event *event_wait_list,
cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueMigrateMemObjects);
return 0;
return command_queue->cast<CommandQueue>().enqueueMigrateMemObjects(
num_mem_objects, mem_objects, flags, num_events_in_wait_list, event_wait_list, event);
}
cl_int EnqueueNDRangeKernel(cl_command_queue command_queue,
......@@ -982,8 +980,9 @@ cl_int EnqueueNDRangeKernel(cl_command_queue command_queue,
const cl_event *event_wait_list,
cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueNDRangeKernel);
return 0;
return command_queue->cast<CommandQueue>().enqueueNDRangeKernel(
kernel, work_dim, global_work_offset, global_work_size, local_work_size,
num_events_in_wait_list, event_wait_list, event);
}
cl_int EnqueueNativeKernel(cl_command_queue command_queue,
......@@ -997,8 +996,9 @@ cl_int EnqueueNativeKernel(cl_command_queue command_queue,
const cl_event *event_wait_list,
cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueNativeKernel);
return 0;
return command_queue->cast<CommandQueue>().enqueueNativeKernel(
user_func, args, cb_args, num_mem_objects, mem_list, args_mem_loc, num_events_in_wait_list,
event_wait_list, event);
}
cl_int EnqueueMarkerWithWaitList(cl_command_queue command_queue,
......@@ -1006,8 +1006,8 @@ cl_int EnqueueMarkerWithWaitList(cl_command_queue command_queue,
const cl_event *event_wait_list,
cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueMarkerWithWaitList);
return 0;
return command_queue->cast<CommandQueue>().enqueueMarkerWithWaitList(num_events_in_wait_list,
event_wait_list, event);
}
cl_int EnqueueBarrierWithWaitList(cl_command_queue command_queue,
......@@ -1015,8 +1015,8 @@ cl_int EnqueueBarrierWithWaitList(cl_command_queue command_queue,
const cl_event *event_wait_list,
cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueBarrierWithWaitList);
return 0;
return command_queue->cast<CommandQueue>().enqueueBarrierWithWaitList(num_events_in_wait_list,
event_wait_list, event);
}
cl_int EnqueueSVMFree(cl_command_queue command_queue,
......@@ -1141,22 +1141,19 @@ cl_mem CreateImage3D(cl_context context,
cl_int EnqueueMarker(cl_command_queue command_queue, cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueMarker);
return 0;
return command_queue->cast<CommandQueue>().enqueueMarker(event);
}
cl_int EnqueueWaitForEvents(cl_command_queue command_queue,
cl_uint num_events,
const cl_event *event_list)
{
WARN_NOT_SUPPORTED(EnqueueWaitForEvents);
return 0;
return command_queue->cast<CommandQueue>().enqueueWaitForEvents(num_events, event_list);
}
cl_int EnqueueBarrier(cl_command_queue command_queue)
{
WARN_NOT_SUPPORTED(EnqueueBarrier);
return 0;
return command_queue->cast<CommandQueue>().enqueueBarrier();
}
cl_int UnloadCompiler()
......@@ -1200,8 +1197,8 @@ cl_int EnqueueTask(cl_command_queue command_queue,
const cl_event *event_wait_list,
cl_event *event)
{
WARN_NOT_SUPPORTED(EnqueueTask);
return 0;
return command_queue->cast<CommandQueue>().enqueueTask(kernel, num_events_in_wait_list,
event_wait_list, event);
}
} // namespace cl
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment