Merge remote-tracking branch 'origin/sycl' into HEAD

Author: aelovikov-intel

clang/lib/Driver/ToolChains/Cuda.cpp

@@ -968,6 +968,11 @@ void CudaToolChain::addClangTargetOptions(
                       "--nvptx-prec-sqrtf32=0"});
 
     CC1Args.append({"-mllvm", "-enable-memcpyopt-without-libcalls"});
+
+    if (DriverArgs.hasFlag(options::OPT_fsycl_id_queries_fit_in_int,
+                           options::OPT_fno_sycl_id_queries_fit_in_int, false))
+      CC1Args.append(
+          {"-mllvm", "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"});
   } else {
     CC1Args.append({"-fcuda-is-device", "-mllvm",
                     "-enable-memcpyopt-without-libcalls",

clang/test/Driver/sycl-nvptx-id-queries-fit-in-int.cpp

@@ -0,0 +1,16 @@
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fno-sycl-id-queries-fit-in-int %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-fit-in-int %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-INT %s
+
+// CHECK-INT: "-mllvm" "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"
+// CHECK-DEFAULT-NOT: "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"

libclc/libspirv/lib/ptx-nvidiacl/workitem/get_global_id.cl

@@ -8,17 +8,31 @@
 
 #include <libspirv/spirv.h>
 
+extern int __nvvm_reflect_ocl(constant char *);
+
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_x() {
+  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT")) {
+    return (uint)__spirv_WorkgroupId_x() * (uint)__spirv_WorkgroupSize_x() +
+           (uint)__spirv_LocalInvocationId_x() + (uint)__spirv_GlobalOffset_x();
+  }
   return __spirv_WorkgroupId_x() * __spirv_WorkgroupSize_x() +
          __spirv_LocalInvocationId_x() + __spirv_GlobalOffset_x();
 }
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_y() {
+  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT")) {
+    return (uint)__spirv_WorkgroupId_y() * (uint)__spirv_WorkgroupSize_y() +
+           (uint)__spirv_LocalInvocationId_y() + (uint)__spirv_GlobalOffset_y();
+  }
   return __spirv_WorkgroupId_y() * __spirv_WorkgroupSize_y() +
          __spirv_LocalInvocationId_y() + __spirv_GlobalOffset_y();
 }
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_z() {
+  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT")) {
+    return (uint)__spirv_WorkgroupId_z() * (uint)__spirv_WorkgroupSize_z() +
+           (uint)__spirv_LocalInvocationId_z() + (uint)__spirv_GlobalOffset_z();
+  }
   return __spirv_WorkgroupId_z() * __spirv_WorkgroupSize_z() +
          __spirv_LocalInvocationId_z() + __spirv_GlobalOffset_z();
 }

libdevice/crt_wrapper.cpp

@@ -12,10 +12,8 @@
 
 #include <cstdint>
 
-#ifndef __NVPTX__
 #define RAND_NEXT_LEN 1024
 DeviceGlobal<uint64_t[RAND_NEXT_LEN]> RandNext;
-#endif
 
 #if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) ||           \
     defined(__AMDGCN__)
@@ -34,8 +32,6 @@ int memcmp(const void *s1, const void *s2, size_t n) {
   return __devicelib_memcmp(s1, s2, n);
 }
 
-#ifndef __NVPTX__
-
 // This simple rand is for ease of use only, the implementation aligns with
 // LLVM libc rand which is based on xorshift64star pseudo random number
 // generator. If work item number <= 1024, each work item has its own internal
@@ -107,8 +103,6 @@ void srand(unsigned int seed) {
   RAND_NEXT_ACC[gid1] = seed;
 }
 
-#endif
-
 #if defined(_WIN32)
 // Truncates a wide (16 or 32 bit) string (wstr) into an ASCII string (str).
 // Any non-ASCII characters are replaced by question mark '?'.

sycl/doc/design/SYCLBINDesign.md

@@ -225,7 +225,7 @@ clang-offload-packager invocation to clang-linker-wrapper together with the new
 `--syclbin` flag.
 
 Setting this option will override `-fsycl`. Passing`-fsycl-device-only` with
-`-fsyclbin` will cause `-fsycl-device-only` to be considered unused.
+`-fsyclbin` will cause `-fsyclbin` to be considered unused.
 
 The behavior is dependent on using the clang-linker-wrapper. As the current
 default offload compilation behavior is using the old offload model (driver

sycl/source/detail/cg.hpp

@@ -725,14 +725,10 @@ class CGHostTask : public CG {
   std::shared_ptr<detail::context_impl> MContext;
   std::vector<ArgDesc> MArgs;
 
-  CGHostTask(std::shared_ptr<HostTask> HostTask,
-             std::shared_ptr<detail::queue_impl> Queue,
+  CGHostTask(std::shared_ptr<HostTask> HostTask, detail::queue_impl *Queue,
              std::shared_ptr<detail::context_impl> Context,
              std::vector<ArgDesc> Args, CG::StorageInitHelper CGData,
-             CGType Type, detail::code_location loc = {})
-      : CG(Type, std::move(CGData), std::move(loc)),
-        MHostTask(std::move(HostTask)), MQueue(Queue), MContext(Context),
-        MArgs(std::move(Args)) {}
+             CGType Type, detail::code_location loc = {});
 };
 
 } // namespace detail

sycl/source/detail/event_impl.cpp

@@ -218,7 +218,7 @@ void event_impl::setQueue(queue_impl &Queue) {
 void event_impl::setSubmittedQueue(std::weak_ptr<queue_impl> SubmittedQueue) {
   MSubmittedQueue = std::move(SubmittedQueue);
   if (MHostProfilingInfo) {
-    if (auto QueuePtr = MSubmittedQueue.lock()) {
+    if (std::shared_ptr<queue_impl> QueuePtr = MSubmittedQueue.lock()) {
       device_impl &Device = QueuePtr->getDeviceImpl();
       MHostProfilingInfo->setDevice(&Device);
     }
@@ -251,7 +251,7 @@ void *event_impl::instrumentationProlog(std::string &Name, int32_t StreamID,
     // queue is available with the wait events. We check to see if the
     // TraceEvent is available in the Queue object.
     void *TraceEvent = nullptr;
-    if (QueueImplPtr Queue = MQueue.lock()) {
+    if (std::shared_ptr<queue_impl> Queue = MQueue.lock()) {
       TraceEvent = Queue->getTraceEvent();
       WaitEvent =
           (TraceEvent ? static_cast<xpti_td *>(TraceEvent) : GSYCLGraphEvent);
@@ -320,7 +320,7 @@ void event_impl::wait_and_throw(
     std::shared_ptr<sycl::detail::event_impl> Self) {
   wait(Self);
 
-  if (QueueImplPtr SubmittedQueue = MSubmittedQueue.lock())
+  if (std::shared_ptr<queue_impl> SubmittedQueue = MSubmittedQueue.lock())
     SubmittedQueue->throw_asynchronous();
 }
 
@@ -465,7 +465,7 @@ event_impl::get_backend_info<info::platform::version>() const {
                           "the info::platform::version info descriptor can "
                           "only be queried with an OpenCL backend");
   }
-  if (QueueImplPtr Queue = MQueue.lock()) {
+  if (std::shared_ptr<queue_impl> Queue = MQueue.lock()) {
     return Queue->getDeviceImpl()
         .get_platform()
         .get_info<info::platform::version>();
@@ -488,7 +488,7 @@ event_impl::get_backend_info<info::device::version>() const {
                           "the info::device::version info descriptor can only "
                           "be queried with an OpenCL backend");
   }
-  if (QueueImplPtr Queue = MQueue.lock()) {
+  if (std::shared_ptr<queue_impl> Queue = MQueue.lock()) {
     return Queue->getDeviceImpl().get_info<info::device::version>();
   }
   return ""; // If the queue has been released, no device will be associated so
@@ -555,21 +555,21 @@ std::vector<EventImplPtr> event_impl::getWaitList() {
   return Result;
 }
 
-void event_impl::flushIfNeeded(const QueueImplPtr &UserQueue) {
+void event_impl::flushIfNeeded(queue_impl *UserQueue) {
   // Some events might not have a native handle underneath even at this point,
   // e.g. those produced by memset with 0 size (no UR call is made).
   auto Handle = this->getHandle();
   if (MIsFlushed || !Handle)
     return;
 
-  QueueImplPtr Queue = MQueue.lock();
+  std::shared_ptr<queue_impl> Queue = MQueue.lock();
   // If the queue has been released, all of the commands have already been
   // implicitly flushed by urQueueRelease.
   if (!Queue) {
     MIsFlushed = true;
     return;
   }
-  if (Queue == UserQueue)
+  if (Queue.get() == UserQueue)
     return;
 
   // Check if the task for this event has already been submitted.
@@ -607,9 +607,9 @@ void event_impl::setSubmissionTime() {
   if (!MIsProfilingEnabled && !MProfilingTagEvent)
     return;
 
-  std::weak_ptr<queue_impl> Queue = isHost() ? MSubmittedQueue : MQueue;
-  if (QueueImplPtr QueuePtr = Queue.lock()) {
-    device_impl &Device = QueuePtr->getDeviceImpl();
+  if (std::shared_ptr<queue_impl> Queue =
+          isHost() ? MSubmittedQueue.lock() : MQueue.lock()) {
+    device_impl &Device = Queue->getDeviceImpl();
     MSubmitTime = getTimestamp(&Device);
   }
 }

sycl/source/detail/event_impl.hpp

@@ -31,7 +31,6 @@ class Adapter;
 class context_impl;
 using ContextImplPtr = std::shared_ptr<sycl::detail::context_impl>;
 class queue_impl;
-using QueueImplPtr = std::shared_ptr<sycl::detail::queue_impl>;
 class event_impl;
 using EventImplPtr = std::shared_ptr<sycl::detail::event_impl>;
 
@@ -238,7 +237,7 @@ class event_impl : public std::enable_shared_from_this<event_impl> {
   /// Performs a flush on the queue associated with this event if the user queue
   /// is different and the task associated with this event hasn't been submitted
   /// to the device yet.
-  void flushIfNeeded(const QueueImplPtr &UserQueue);
+  void flushIfNeeded(queue_impl *UserQueue);
 
   /// Cleans dependencies of this event_impl.
   void cleanupDependencyEvents();
@@ -258,7 +257,9 @@ class event_impl : public std::enable_shared_from_this<event_impl> {
   ///
   /// @return shared_ptr to MWorkerQueue, please be aware it can be empty
   /// pointer
-  QueueImplPtr getWorkerQueue() { return MWorkerQueue.lock(); };
+  std::shared_ptr<sycl::detail::queue_impl> getWorkerQueue() {
+    return MWorkerQueue.lock();
+  };
 
   /// Sets worker queue for command.
   ///
@@ -285,7 +286,9 @@ class event_impl : public std::enable_shared_from_this<event_impl> {
   /// @return Submission time for command associated with this event
   uint64_t getSubmissionTime();
 
-  QueueImplPtr getSubmittedQueue() const { return MSubmittedQueue.lock(); };
+  std::shared_ptr<sycl::detail::queue_impl> getSubmittedQueue() const {
+    return MSubmittedQueue.lock();
+  };
 
   /// Checks if this event is complete.
   ///

sycl/source/detail/graph_impl.hpp

@@ -288,7 +288,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
 
       return std::make_unique<sycl::detail::CGHostTask>(
           sycl::detail::CGHostTask(
-              std::move(HostTaskSPtr), CommandGroupPtr->MQueue,
+              std::move(HostTaskSPtr), CommandGroupPtr->MQueue.get(),
               CommandGroupPtr->MContext, std::move(NewArgs), std::move(Data),
               CommandGroupPtr->getType(), Loc));
     }

sycl/source/detail/queue_impl.hpp

@@ -650,9 +650,12 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
   // for in order ones.
   void revisitUnenqueuedCommandsState(const EventImplPtr &CompletedHostTask);
 
-  static ContextImplPtr getContext(const QueueImplPtr &Queue) {
+  static ContextImplPtr getContext(queue_impl *Queue) {
     return Queue ? Queue->getContextImplPtr() : nullptr;
   }
+  static ContextImplPtr getContext(const QueueImplPtr &Queue) {
+    return getContext(Queue.get());
+  }
 
   // Must be called under MMutex protection
   void doUnenqueuedCommandCleanup(