[SYCL] Avoid using intrinsics to get global id in stream (#2318)

Avoid calling intrinsics to get global id when calculating the offset of
the work item in the flush buffer. When stream is used in a single_task
kernel this could cause an overhead on targets like FPGA. That is why
use global atomic variable to calculate offsets.

Signed-off-by: Artur Gainullin <[email protected]>

Author: againull

sycl/include/CL/sycl/stream.hpp

@@ -811,13 +811,12 @@ class __SYCL_EXPORT stream {
 
 #ifdef __SYCL_DEVICE_ONLY__
   void __init() {
-    // Calculate work item's global id, this should be done once, that
-    // is why this is done in _init method, call to __init method is generated
-    // by frontend. As a result each work item will write to its own section
-    // in the flush buffer
-
-    id<1> GlobalId = __spirv::initGlobalInvocationId<1, id<1>>();
-    WIOffset = (unsigned int)GlobalId[0] * FlushBufferSize;
+    // Calculate offset in the flush buffer for each work item in the global
+    // work space. We need to avoid calling intrinsics to get global id because
+    // when stream is used in  a single_task kernel this could cause some
+    // overhead on FPGA target. That is why use global atomic variable to
+    // calculate offsets.
+    WIOffset = GlobalOffset[1].fetch_add(FlushBufferSize);
   }
 
   void __finalize() {

sycl/source/detail/stream_impl.hpp

@@ -38,12 +38,12 @@ class __SYCL_EXPORT stream_impl {
   }
 
   // Method to provide an atomic access to the offset in the global stream
-  // buffer
+  // buffer and offset in the flush buffer
   GlobalOffsetAccessorT accessGlobalOffset(handler &CGH) {
     auto OffsetSubBuf = buffer<char, 1>(Buf, id<1>(0), range<1>(OffsetSize));
-    auto ReinterpretedBuf = OffsetSubBuf.reinterpret<unsigned, 1>(range<1>(1));
+    auto ReinterpretedBuf = OffsetSubBuf.reinterpret<unsigned, 1>(range<1>(2));
     return ReinterpretedBuf.get_access<cl::sycl::access::mode::atomic>(
-        CGH, range<1>(1), id<1>(0));
+        CGH, range<1>(2), id<1>(0));
   }
 
   // Copy stream buffer to the host and print the contents
@@ -61,10 +61,9 @@ class __SYCL_EXPORT stream_impl {
   // statement till the semicolon
   unsigned MaxStatementSize_;
 
-  // Size of the variable which is used as an offset in the stream buffer.
   // Additinonal memory is allocated in the beginning of the stream buffer for
-  // this variable.
-  static const size_t OffsetSize = sizeof(unsigned);
+  // 2 variables: offset in the stream buffer and offset in the flush buffer.
+  static const size_t OffsetSize = 2 * sizeof(unsigned);
 
   // Vector on the host side which is used to initialize the stream buffer
   std::vector<char> Data;

sycl/test/basic_tests/stream/no_intrinsic.cpp

@@ -0,0 +1,32 @@
+// RUN: %clangxx -S -emit-llvm -fsycl-device-only %s -o - -Xclang -disable-llvm-passes | FileCheck %s
+
+//==------------------ no_intrinsics.cpp - SYCL stream 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
+//
+//===----------------------------------------------------------------------===//
+
+// Test to check that intrinsics to get a global id are not generated for the
+// stream.
+
+// CHECK-NOT: call spir_func void @{{.*}}__spirvL22initGlobalInvocationId{{.*}}
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+int main() {
+  {
+    queue Queue;
+
+    Queue.submit([&](handler &CGH) {
+      stream Out(1024, 80, CGH);
+      CGH.single_task<class integral>([=]() { Out << "Hello, World!\n"; });
+    });
+    Queue.wait();
+  }
+
+  return 0;
+}