[SYCL] Add inline to get_local_linear_id (#6302)

* [SYCL] Add inline to get_local_linear_id

These changes make sycl::detail::get_local_linear_id and its
specializations inline.

Signed-off-by: Larsen, Steffen <[email protected]>
Co-authored-by: Sergey Semenov <[email protected]>

Author: steffenlarsen

sycl/include/CL/sycl/group_algorithm.hpp

@@ -65,12 +65,12 @@ get_local_linear_range<ext::oneapi::sub_group>(ext::oneapi::sub_group g) {
 
 // ---- get_local_linear_id
 template <typename Group>
-typename Group::linear_id_type get_local_linear_id(Group g);
+inline typename Group::linear_id_type get_local_linear_id(Group g);
 
 #ifdef __SYCL_DEVICE_ONLY__
 #define __SYCL_GROUP_GET_LOCAL_LINEAR_ID(D)                                    \
   template <>                                                                  \
-  group<D>::linear_id_type get_local_linear_id<group<D>>(group<D>) {           \
+  inline group<D>::linear_id_type get_local_linear_id<group<D>>(group<D>) {    \
     nd_item<D> it = cl::sycl::detail::Builder::getNDItem<D>();                 \
     return it.get_local_linear_id();                                           \
   }

sycl/test/regression/joint_group_conflict.cpp

@@ -0,0 +1,25 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fPIC -DCASE1 %s -c -o %t.1.o
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fPIC -DCASE2 %s -c -o %t.2.o
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -shared %t.1.o %t.2.o -o %t.so
+
+// Tests that creating a shared library with multiple object files using joint
+// group operations does not cause conflicting definitions.
+
+#include <CL/sycl.hpp>
+
+#ifdef CASE1
+#define FNAME test1
+#define KNAME kernel1
+#else
+#define FNAME test2
+#define KNAME kernel2
+#endif
+
+void FNAME(sycl::queue &Q, bool *In, bool *Out) {
+  sycl::nd_range<1> WorkRange(sycl::range<1>(8), sycl::range<1>(16));
+  Q.parallel_for<class KNAME>(WorkRange, [=](sycl::nd_item<1> It) {
+     sycl::group<1> Group = It.get_group();
+     size_t I = It.get_global_linear_id();
+     Out[I] = sycl::joint_any_of(Group, In, In, [](bool B) { return B; });
+   }).wait();
+}