Add clEnqueueNDRangeKernelINTEL API

Related-To: NEO-2712

Change-Id: If1d16d9d626871a9dc4b19282f9edc5786ffa398
Signed-off-by: Filip Hazubski <filip.hazubski@intel.com>

Author: fhazubski-Intel

core/program/CMakeLists.txt

@@ -8,6 +8,8 @@ set(NEO_CORE_PROGRAM
   ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
   ${CMAKE_CURRENT_SOURCE_DIR}/print_formatter.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/print_formatter.h
+  ${CMAKE_CURRENT_SOURCE_DIR}/sync_buffer_handler.cpp
+  ${CMAKE_CURRENT_SOURCE_DIR}/sync_buffer_handler.h
 )
 
 set_property(GLOBAL PROPERTY NEO_CORE_PROGRAM ${NEO_CORE_PROGRAM})

core/program/sync_buffer_handler.cpp

@@ -0,0 +1,54 @@
+/*
+ * Copyright (C) 2019 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "core/program/sync_buffer_handler.h"
+
+#include "core/memory_manager/graphics_allocation.h"
+#include "runtime/command_stream/command_stream_receiver.h"
+#include "runtime/kernel/kernel.h"
+#include "runtime/memory_manager/memory_manager.h"
+
+namespace NEO {
+
+SyncBufferHandler::~SyncBufferHandler() {
+    memoryManager.checkGpuUsageAndDestroyGraphicsAllocations(graphicsAllocation);
+};
+SyncBufferHandler::SyncBufferHandler(Device &device)
+    : device(device), memoryManager(*device.getMemoryManager()) {
+
+    allocateNewBuffer();
+}
+
+void SyncBufferHandler::prepareForEnqueue(size_t workGroupsCount, Kernel &kernel, CommandStreamReceiver &csr) {
+    auto requiredSize = workGroupsCount;
+    std::lock_guard<std::mutex> guard(this->mutex);
+
+    bool isCurrentBufferFull = (usedBufferSize + requiredSize > bufferSize);
+    if (isCurrentBufferFull) {
+        memoryManager.checkGpuUsageAndDestroyGraphicsAllocations(graphicsAllocation);
+        allocateNewBuffer();
+        usedBufferSize = 0;
+    }
+
+    kernel.patchSyncBuffer(device, graphicsAllocation, usedBufferSize);
+    csr.makeResident(*graphicsAllocation);
+
+    usedBufferSize += requiredSize;
+}
+
+void SyncBufferHandler::allocateNewBuffer() {
+    AllocationProperties allocationProperties{device.getRootDeviceIndex(), true, bufferSize,
+                                              GraphicsAllocation::AllocationType::LINEAR_STREAM,
+                                              false, false, static_cast<uint32_t>(device.getDeviceBitfield().to_ulong())};
+    graphicsAllocation = memoryManager.allocateGraphicsMemoryWithProperties(allocationProperties);
+    UNRECOVERABLE_IF(graphicsAllocation == nullptr);
+
+    auto cpuPointer = graphicsAllocation->getUnderlyingBuffer();
+    std::memset(cpuPointer, 0, bufferSize);
+}
+
+} // namespace NEO

core/program/sync_buffer_handler.h

@@ -0,0 +1,42 @@
+/*
+ * Copyright (C) 2019 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#pragma once
+
+#include "core/helpers/basic_math.h"
+
+#include <mutex>
+
+namespace NEO {
+
+class CommandStreamReceiver;
+class Context;
+class Device;
+class GraphicsAllocation;
+class MemoryManager;
+class Kernel;
+
+class SyncBufferHandler {
+  public:
+    ~SyncBufferHandler();
+
+    SyncBufferHandler(Device &device);
+
+    void prepareForEnqueue(size_t workGroupsCount, Kernel &kernel, CommandStreamReceiver &csr);
+
+  protected:
+    void allocateNewBuffer();
+
+    Device &device;
+    MemoryManager &memoryManager;
+    GraphicsAllocation *graphicsAllocation;
+    const size_t bufferSize = 64 * KB;
+    size_t usedBufferSize = 0;
+    std::mutex mutex;
+};
+
+} // namespace NEO

runtime/api/api.cpp

@@ -3135,18 +3135,24 @@ cl_int CL_API_CALL clEnqueueNDRangeKernel(cl_command_queue commandQueue,
                    "event", DebugManager.getEvents(reinterpret_cast<const uintptr_t *>(event), 1));
 
     CommandQueue *pCommandQueue = nullptr;
+    Kernel *pKernel = nullptr;
 
     retVal = validateObjects(
         WithCastToInternal(commandQueue, &pCommandQueue),
-        kernel,
+        WithCastToInternal(kernel, &pKernel),
         EventWaitList(numEventsInWaitList, eventWaitList));
 
     if (CL_SUCCESS != retVal) {
         TRACING_EXIT(clEnqueueNDRangeKernel, &retVal);
         return retVal;
     }
 
-    auto pKernel = castToObjectOrAbort<Kernel>(kernel);
+    if (pKernel->getKernelInfo().patchInfo.pAllocateSyncBuffer != nullptr) {
+        retVal = CL_INVALID_KERNEL;
+        TRACING_EXIT(clEnqueueNDRangeKernel, &retVal);
+        return retVal;
+    }
+
     TakeOwnershipWrapper<Kernel> kernelOwnership(*pKernel, gtpinIsGTPinInitialized());
     if (gtpinIsGTPinInitialized()) {
         gtpinNotifyKernelSubmit(kernel, pCommandQueue);
@@ -3947,6 +3953,7 @@ void *CL_API_CALL clGetExtensionFunctionAddress(const char *funcName) {
     RETURN_FUNC_PTR_IF_EXIST(clGetDeviceFunctionPointerINTEL);
     RETURN_FUNC_PTR_IF_EXIST(clGetDeviceGlobalVariablePointerINTEL);
     RETURN_FUNC_PTR_IF_EXIST(clGetExecutionInfoINTEL);
+    RETURN_FUNC_PTR_IF_EXIST(clEnqueueNDRangeKernelINTEL);
 
     void *ret = sharingFactory.getExtensionFunctionAddress(funcName);
     if (ret != nullptr) {
@@ -5197,3 +5204,70 @@ cl_int CL_API_CALL clGetExecutionInfoINTEL(cl_command_queue commandQueue,
 
     return retVal;
 }
+
+cl_int CL_API_CALL clEnqueueNDRangeKernelINTEL(cl_command_queue commandQueue,
+                                               cl_kernel kernel,
+                                               cl_uint workDim,
+                                               const size_t *globalWorkOffset,
+                                               const size_t *workgroupCount,
+                                               const size_t *localWorkSize,
+                                               cl_uint numEventsInWaitList,
+                                               const cl_event *eventWaitList,
+                                               cl_event *event) {
+    cl_int retVal = CL_SUCCESS;
+    API_ENTER(&retVal);
+    DBG_LOG_INPUTS("commandQueue", commandQueue, "cl_kernel", kernel,
+                   "globalWorkOffset[0]", DebugManager.getInput(globalWorkOffset, 0),
+                   "globalWorkOffset[1]", DebugManager.getInput(globalWorkOffset, 1),
+                   "globalWorkOffset[2]", DebugManager.getInput(globalWorkOffset, 2),
+                   "workgroupCount", DebugManager.getSizes(workgroupCount, workDim, false),
+                   "localWorkSize", DebugManager.getSizes(localWorkSize, workDim, true),
+                   "numEventsInWaitList", numEventsInWaitList,
+                   "eventWaitList", DebugManager.getEvents(reinterpret_cast<const uintptr_t *>(eventWaitList), numEventsInWaitList),
+                   "event", DebugManager.getEvents(reinterpret_cast<const uintptr_t *>(event), 1));
+
+    CommandQueue *pCommandQueue = nullptr;
+    Kernel *pKernel = nullptr;
+
+    retVal = validateObjects(
+        WithCastToInternal(commandQueue, &pCommandQueue),
+        WithCastToInternal(kernel, &pKernel),
+        EventWaitList(numEventsInWaitList, eventWaitList));
+
+    if (CL_SUCCESS != retVal) {
+        return retVal;
+    }
+
+    size_t globalWorkSize[3];
+    size_t requestedNumberOfWorkgroups = 1;
+    for (size_t i = 0; i < workDim; i++) {
+        globalWorkSize[i] = workgroupCount[i] * localWorkSize[i];
+        requestedNumberOfWorkgroups *= workgroupCount[i];
+    }
+
+    size_t maximalNumberOfWorkgroupsAllowed = pKernel->getMaxWorkGroupCount(workDim, localWorkSize);
+    if (requestedNumberOfWorkgroups > maximalNumberOfWorkgroupsAllowed) {
+        retVal = CL_INVALID_VALUE;
+        return retVal;
+    }
+
+    TakeOwnershipWrapper<Kernel> kernelOwnership(*pKernel, gtpinIsGTPinInitialized());
+    if (gtpinIsGTPinInitialized()) {
+        gtpinNotifyKernelSubmit(kernel, pCommandQueue);
+    }
+
+    pCommandQueue->getDevice().allocateSyncBufferHandler();
+
+    retVal = pCommandQueue->enqueueKernel(
+        kernel,
+        workDim,
+        globalWorkOffset,
+        globalWorkSize,
+        localWorkSize,
+        numEventsInWaitList,
+        eventWaitList,
+        event);
+
+    DBG_LOG_INPUTS("event", DebugManager.getEvents(reinterpret_cast<const uintptr_t *>(event), 1u));
+    return retVal;
+}

runtime/api/api.h

@@ -1036,6 +1036,16 @@ cl_int CL_API_CALL clGetExecutionInfoINTEL(
     void *paramValue,
     size_t *paramValueSizeRet);
 
+cl_int CL_API_CALL clEnqueueNDRangeKernelINTEL(cl_command_queue commandQueue,
+                                               cl_kernel kernel,
+                                               cl_uint workDim,
+                                               const size_t *globalWorkOffset,
+                                               const size_t *workgroupCount,
+                                               const size_t *localWorkSize,
+                                               cl_uint numEventsInWaitList,
+                                               const cl_event *eventWaitList,
+                                               cl_event *event);
+
 // OpenCL 2.2
 
 cl_int CL_API_CALL clSetProgramSpecializationConstant(

runtime/command_queue/enqueue_common.h

@@ -6,6 +6,7 @@
  */
 
 #pragma once
+#include "core/program/sync_buffer_handler.h"
 #include "core/utilities/range.h"
 #include "runtime/built_ins/built_ins.h"
 #include "runtime/built_ins/builtins_dispatch_builder.h"
@@ -642,6 +643,15 @@ CompletionStamp CommandQueueHw<GfxFamily>::enqueueNonBlocked(
         blocking = true;
         printfHandler->makeResident(getGpgpuCommandStreamReceiver());
     }
+
+    if (multiDispatchInfo.peekMainKernel()->usesSyncBuffer()) {
+        auto &gws = multiDispatchInfo.begin()->getGWS();
+        auto &lws = multiDispatchInfo.begin()->getLocalWorkgroupSize();
+        size_t workGroupsCount = (gws.x * gws.y * gws.z) /
+                                 (lws.x * lws.y * lws.z);
+        device->syncBufferHandler->prepareForEnqueue(workGroupsCount, *multiDispatchInfo.peekMainKernel(), getGpgpuCommandStreamReceiver());
+    }
+
     if (timestampPacketContainer) {
         timestampPacketContainer->makeResident(getGpgpuCommandStreamReceiver());
         timestampPacketDependencies.previousEnqueueNodes.makeResident(getGpgpuCommandStreamReceiver());

runtime/compiler_interface/patchtokens_decoder.cpp

@@ -407,6 +407,10 @@ inline bool decodeToken(const SPatchItemHeader *token, KernelFromPatchtokens &ou
         auto tokDataP = reinterpret_cast<const SPatchDataParameterBuffer *>(token);
         decodeKernelDataParameterToken(tokDataP, out);
     } break;
+
+    case PATCH_TOKEN_ALLOCATE_SYNC_BUFFER: {
+        assignToken(out.tokens.allocateSyncBuffer, token);
+    } break;
     }
 
     return out.decodeStatus != DecoderError::InvalidBinary;

runtime/compiler_interface/patchtokens_decoder.h

@@ -129,6 +129,7 @@ struct KernelFromPatchtokens {
         const SPatchAllocateStatelessPrintfSurface *allocateStatelessPrintfSurface = nullptr;
         const SPatchAllocateStatelessEventPoolSurface *allocateStatelessEventPoolSurface = nullptr;
         const SPatchAllocateStatelessDefaultDeviceQueueSurface *allocateStatelessDefaultDeviceQueueSurface = nullptr;
+        const SPatchAllocateSyncBuffer *allocateSyncBuffer = nullptr;
         const SPatchItemHeader *inlineVmeSamplerInfo = nullptr;
         const SPatchGtpinFreeGRFInfo *gtpinFreeGrfInfo = nullptr;
         const SPatchStateSIP *stateSip = nullptr;

runtime/device/device.cpp

@@ -9,6 +9,7 @@
 
 #include "core/command_stream/preemption.h"
 #include "core/helpers/hw_helper.h"
+#include "core/program/sync_buffer_handler.h"
 #include "runtime/command_stream/command_stream_receiver.h"
 #include "runtime/command_stream/experimental_command_buffer.h"
 #include "runtime/device/device_vector.h"
@@ -60,6 +61,7 @@ Device::Device(ExecutionEnvironment *executionEnvironment)
 
 Device::~Device() {
     DEBUG_BREAK_IF(nullptr == executionEnvironment->memoryManager.get());
+    syncBufferHandler.reset();
     if (performanceCounters) {
         performanceCounters->shutdown();
     }
@@ -203,6 +205,15 @@ double Device::getPlatformHostTimerResolution() const {
         return osTime->getHostTimerResolution();
     return 0.0;
 }
+
+void Device::allocateSyncBufferHandler() {
+    TakeOwnershipWrapper<Device> lock(*this);
+    if (syncBufferHandler.get() == nullptr) {
+        syncBufferHandler = std::make_unique<SyncBufferHandler>(*this);
+        UNRECOVERABLE_IF(syncBufferHandler.get() == nullptr);
+    }
+}
+
 GFXCORE_FAMILY Device::getRenderCoreFamily() const {
     return this->getHardwareInfo().platform.eRenderCoreFamily;
 }

runtime/device/device.h

@@ -18,8 +18,9 @@
 #include "engine_node.h"
 
 namespace NEO {
-class OSTime;
 class DriverInfo;
+class OSTime;
+class SyncBufferHandler;
 
 template <>
 struct OpenCLObjectMapper<_cl_device_id> {
@@ -72,6 +73,7 @@ class Device : public BaseObject<_cl_device_id> {
     double getPlatformHostTimerResolution() const;
     bool isSimulation() const;
     GFXCORE_FAMILY getRenderCoreFamily() const;
+    void allocateSyncBufferHandler();
     PerformanceCounters *getPerformanceCounters() { return performanceCounters.get(); }
     PreemptionMode getPreemptionMode() const { return preemptionMode; }
     MOCKABLE_VIRTUAL bool isSourceLevelDebuggerActive() const;
@@ -88,8 +90,10 @@ class Device : public BaseObject<_cl_device_id> {
     virtual uint32_t getRootDeviceIndex() const = 0;
     virtual uint32_t getNumAvailableDevices() const = 0;
     virtual Device *getDeviceById(uint32_t deviceId) const = 0;
+    virtual DeviceBitfield getDeviceBitfield() const = 0;
 
     static decltype(&PerformanceCounters::create) createPerformanceCountersFunc;
+    std::unique_ptr<SyncBufferHandler> syncBufferHandler;
 
   protected:
     Device() = delete;
@@ -113,8 +117,6 @@ class Device : public BaseObject<_cl_device_id> {
     bool createEngine(uint32_t deviceCsrIndex, aub_stream::EngineType engineType);
     MOCKABLE_VIRTUAL std::unique_ptr<CommandStreamReceiver> createCommandStreamReceiver() const;
 
-    virtual DeviceBitfield getDeviceBitfield() const = 0;
-
     std::vector<unsigned int> simultaneousInterops;
     unsigned int enabledClVersion = 0u;
     std::string deviceExtensions;