[SYCL][USM] Enable mem_advise APIs (#544)

* Implement mem_advise, add tests, update documentation.

Signed-off-by: James Brodman <james.brodman@intel.com>

Author: jbrodman

sycl/doc/extensions/USM/USM.adoc

@@ -374,22 +374,28 @@ Return value:: none
 
 '''
 ==== Concurrent USM
-Concurrent USM contains all the utility functions of Explicit USM and Restricted USM.  It introduces a new function, `sycl::mem_advise`, that allows programmers to provide additional information to the underlying runtime about how different allocations are used.  
+Concurrent USM contains all the utility functions of Explicit USM and Restricted USM.  It introduces a new function, `sycl::queue::mem_advise`, that allows programmers to provide additional information to the underlying runtime about how different allocations are used.  
 
 ===== Performance Hints
 ===== prefetch
 In Concurrent USM, prefetch commands may be overlapped with kernel execution.
+
 ===== mem_advise
 [source,cpp]
 ----
-void sycl::mem_advise(void *addr, size_t length, int advice);
+class queue {
+ ...
+ public:
+  ...
+  event mem_advise(void *addr, size_t length, int advice);
+};
 ----
 
 Parameters::
  * `void* addr` - address of allocation
  * `size_t length` - number of bytes in the allocation
  * `int advice` - device-defined advice for the specified allocation
-Return Value:: none
+Return Value:: Returns an event representing the operation.
 
 '''
 ==== General

sycl/include/CL/sycl/detail/queue_impl.hpp

@@ -201,8 +201,9 @@ class queue_impl {
     return m_PropList.get_property<propertyT>();
   }
 
-  event memset(void* ptr, int value, size_t count);
-  event memcpy(void* dest, const void* src, size_t count);
+  event memset(void* Ptr, int Value, size_t Count);
+  event memcpy(void* Dest, const void* Src, size_t Count);
+  event mem_advise(const void *Ptr, size_t Length, int Advice);
 
 private:
   template <typename T>

sycl/include/CL/sycl/detail/usm_dispatch.hpp

@@ -43,13 +43,11 @@ class USMDispatcher {
                               cl_mem_migration_flags Flags,
                               pi_uint32 NumEventsInWaitList,
                               const pi_event *EventWaitList, pi_event *Event);
-  pi_result enqueueMemAdvise(pi_queue Queue, void *Ptr, size_t Size,
-                             cl_mem_advice_intel Advice,
-                             pi_uint32 NumEventsInWaitList,
-                             const pi_event *EventWaitList, pi_event *Event);
   pi_result getMemAllocInfo(pi_context Context, const void *Ptr,
                             cl_mem_info_intel ParamName, size_t ParamValueSize,
                             void *ParamValue, size_t *ParamValueSizeRet);
+  void memAdvise(pi_queue Queue, const void *Ptr, size_t Length, int Advice,
+                 pi_event *Event);
 
 private:
   bool mEmulated = false;

sycl/include/CL/sycl/queue.hpp

@@ -104,12 +104,16 @@ class queue {
     return impl->get_property<propertyT>();
   }
 
-  event memset(void* ptr, int value, size_t count) {
-    return impl->memset(ptr, value, count);
+  event memset(void* Ptr, int Value, size_t Count) {
+    return impl->memset(Ptr, Value, Count);
   }
 
-  event memcpy(void* dest, const void* src, size_t count) {
-    return impl->memcpy(dest, src, count);
+  event memcpy(void* Dest, const void* Src, size_t Count) {
+    return impl->memcpy(Dest, Src, Count);
+  }
+
+  event mem_advise(const void *Ptr, size_t Length, int Advice) {
+    return impl->mem_advise(Ptr, Length, Advice);
   }
 
 private:

sycl/source/detail/queue_impl.cpp

@@ -59,6 +59,18 @@ event queue_impl::memcpy(void *Dest, const void *Src, size_t Count) {
 
   return event(Event, Context);
 }
+
+event queue_impl::mem_advise(const void *Ptr, size_t Length, int Advice) {
+  context Context = get_context();
+  std::shared_ptr<usm::USMDispatcher> USMDispatch =
+    getSyclObjImpl(Context)->getUSMDispatch();
+  cl_event Event;
+
+  USMDispatch->memAdvise(getHandleRef(), Ptr, Length, Advice,
+                         reinterpret_cast<pi_event *>(&Event));
+
+  return event(Event, Context);
+}
 } // namespace detail
 } // namespace sycl
 } // namespace cl

sycl/source/detail/usm/usm_dispatch.cpp

@@ -278,37 +278,6 @@ pi_result USMDispatcher::enqueueMigrateMem(pi_queue Queue, const void *Ptr,
   return RetVal;
 }
 
-pi_result USMDispatcher::enqueueMemAdvise(pi_queue Queue, void *Ptr,
-                                          size_t Size,
-                                          cl_mem_advice_intel Advice,
-                                          pi_uint32 NumEventsInWaitList,
-                                          const pi_event *EventWaitList,
-                                          pi_event *Event) {
-  pi_result RetVal = PI_INVALID_OPERATION;
-
-  if (pi::useBackend(pi::Backend::SYCL_BE_PI_OPENCL)) {
-    cl_command_queue CLQueue = pi::cast<cl_command_queue>(Queue);
-
-    if (mEmulated) {
-      // TODO: What should we do here?
-      // This isn't really supported yet.
-      // Advice is typically safe to ignore,
-      //  so a NOP will do.
-      RetVal = pi::cast<pi_result>(clEnqueueMarkerWithWaitList(
-          CLQueue, NumEventsInWaitList,
-          reinterpret_cast<const cl_event *>(EventWaitList),
-          reinterpret_cast<cl_event *>(Event)));
-    } else {
-      RetVal = pi::cast<pi_result>(pfn_clEnqueueMemAdviseINTEL(
-          CLQueue, Ptr, Size, Advice, NumEventsInWaitList,
-          reinterpret_cast<const cl_event *>(EventWaitList),
-          reinterpret_cast<cl_event *>(Event)));
-    }
-  }
-
-  return RetVal;
-}
-
 pi_result USMDispatcher::getMemAllocInfo(pi_context Context, const void *Ptr,
                                          cl_mem_info_intel ParamName,
                                          size_t ParamValueSize,
@@ -337,6 +306,30 @@ pi_result USMDispatcher::getMemAllocInfo(pi_context Context, const void *Ptr,
   return RetVal;
 }
 
+void USMDispatcher::memAdvise(pi_queue Queue, const void *Ptr, size_t Length,
+                              int Advice, pi_event *Event) {
+  if (pi::useBackend(pi::Backend::SYCL_BE_PI_OPENCL)) {
+    cl_command_queue CLQueue = pi::cast<cl_command_queue>(Queue);
+
+    if (mEmulated) {
+      // memAdvise does nothing here
+      PI_CHECK(clEnqueueMarkerWithWaitList(
+          CLQueue, 0, nullptr, reinterpret_cast<cl_event *>(Event)));
+    } else {
+      // Temporary until driver supports
+      // memAdvise doesn't do anything on an iGPU anyway
+      PI_CHECK(clEnqueueMarkerWithWaitList(
+                 CLQueue, 0, nullptr, reinterpret_cast<cl_event *>(Event)));
+      /*
+      // Enable once this is supported in the driver
+      auto CLAdvice = *reinterpret_cast<cl_mem_advice_intel *>(&Advice);
+      PI_CHECK(pfn_clEnqueueMemAdviseINTEL(
+          CLQueue, Ptr, Length, CLAdvice, 0, nullptr,
+          reinterpret_cast<cl_event *>(Event)));
+      */
+    }
+  }
+}
 } // namespace usm
 } // namespace detail
 } // namespace sycl

sycl/test/usm/memadvise.cpp

@@ -0,0 +1,83 @@
+// RUN: %clangxx -fsycl %s -o %t1.out -lOpenCL
+// RUN: %CPU_RUN_PLACEHOLDER %t1.out
+// TODO: SYCL specific fail - analyze and enable
+// XFAIL: windows
+
+//==---------------- memadvise.cpp - Shared Memory Linked List test --------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+int numNodes = 4;
+
+struct Node {
+  Node() : pNext(nullptr), Num(0xDEADBEEF) {}
+
+  Node *pNext;
+  uint32_t Num;
+};
+
+class foo;
+int main() {
+  queue q;
+  auto dev = q.get_device();
+  auto ctxt = q.get_context();
+  Node *s_head = nullptr;
+  Node *s_cur = nullptr;
+
+  s_head = (Node *)malloc_shared(sizeof(Node), dev, ctxt);
+  if (s_head == nullptr) {
+    return -1;
+  }
+  q.mem_advise(s_head, sizeof(Node), 42);
+  s_cur = s_head;
+
+  for (int i = 0; i < numNodes; i++) {
+    s_cur->Num = i * 2;
+
+    if (i != (numNodes - 1)) {
+      s_cur->pNext = (Node *)malloc_shared(sizeof(Node), dev, ctxt);
+      if (s_cur->pNext == nullptr) {
+        return -1;
+      }
+      q.mem_advise(s_cur->pNext, sizeof(Node), 42);
+    } else {
+      s_cur->pNext = nullptr;
+    }
+
+    s_cur = s_cur->pNext;
+  }
+
+  auto e1 = q.submit([=](handler &cgh) {
+    cgh.single_task<class foo>([=]() {
+      Node *pHead = s_head;
+      while (pHead) {
+        pHead->Num = pHead->Num * 2 + 1;
+        pHead = pHead->pNext;
+      }
+    });
+  });
+
+  e1.wait();
+
+  s_cur = s_head;
+  int mismatches = 0;
+  for (int i = 0; i < numNodes; i++) {
+    const int want = i * 4 + 1;
+    if (s_cur->Num != want) {
+      return -1;
+    }
+    Node *old = s_cur;
+    s_cur = s_cur->pNext;
+    free(old, ctxt);
+  }
+
+  return 0;
+}