Multiple definition of `void cl::sycl::group_barrier<...>`

Hi,
compiling and linking multiple source files with compute++ 2.2.1 I get linker errors about multiple definition of void cl::sycl::group_barrier<...>:

/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ScopedContext.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::experimental::sub_group>(cl::sycl::experimental::sub_group, cl::sycl::memory_scope)':
ScopedContext.cc:(.text+0x300): multiple definition of `void cl::sycl::group_barrier<cl::sycl::experimental::sub_group>(cl::sycl::experimental::sub_group, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x300): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ScopedContext.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<1> >(cl::sycl::group<1>, cl::sycl::memory_scope)':
ScopedContext.cc:(.text+0x0): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<1> >(cl::sycl::group<1>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x0): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ScopedContext.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<2> >(cl::sycl::group<2>, cl::sycl::memory_scope)':
ScopedContext.cc:(.text+0x100): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<2> >(cl::sycl::group<2>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x100): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ScopedContext.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<3> >(cl::sycl::group<3>, cl::sycl::memory_scope)':
ScopedContext.cc:(.text+0x200): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<3> >(cl::sycl::group<3>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x200): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ProductBase.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::experimental::sub_group>(cl::sycl::experimental::sub_group, cl::sycl::memory_scope)':
ProductBase.cc:(.text+0x300): multiple definition of `void cl::sycl::group_barrier<cl::sycl::experimental::sub_group>(cl::sycl::experimental::sub_group, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x300): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ProductBase.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<1> >(cl::sycl::group<1>, cl::sycl::memory_scope)':
ProductBase.cc:(.text+0x0): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<1> >(cl::sycl::group<1>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x0): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ProductBase.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<2> >(cl::sycl::group<2>, cl::sycl::memory_scope)':
ProductBase.cc:(.text+0x100): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<2> >(cl::sycl::group<2>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x100): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ProductBase.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<3> >(cl::sycl::group<3>, cl::sycl::memory_scope)':
ProductBase.cc:(.text+0x200): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<3> >(cl::sycl::group<3>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x200): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ContextState.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::experimental::sub_group>(cl::sycl::experimental::sub_group, cl::sycl::memory_scope)':
ContextState.cc:(.text+0x300): multiple definition of `void cl::sycl::group_barrier<cl::sycl::experimental::sub_group>(cl::sycl::experimental::sub_group, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x300): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ContextState.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<1> >(cl::sycl::group<1>, cl::sycl::memory_scope)':
ContextState.cc:(.text+0x0): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<1> >(cl::sycl::group<1>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x0): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ContextState.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<2> >(cl::sycl::group<2>, cl::sycl::memory_scope)':
ContextState.cc:(.text+0x100): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<2> >(cl::sycl::group<2>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x100): first defined here
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/ContextState.cc.o: In function `void cl::sycl::group_barrier<cl::sycl::group<3> >(cl::sycl::group<3>, cl::sycl::memory_scope)':
ContextState.cc:(.text+0x200): multiple definition of `void cl::sycl::group_barrier<cl::sycl::group<3> >(cl::sycl::group<3>, cl::sycl::memory_scope)'
/data/user/fwyzard/pixeltrack-standalone/obj/sycltest/SYCLCore/chooseDevice.cc.o:chooseDevice.cc:(.text+0x200): first defined here
compute++: error: linker command failed with exit code 1 (use -v to see invocation)

The solution I found is to mark the cl::sycl::group_barrier template spacialisations as inline:

diff --git a/SYCL/group_functions.h b/SYCL/group_functions.h
index 4053834..be96643 100644
--- a/SYCL/group_functions.h
+++ b/SYCL/group_functions.h
@@ -31,6 +31,7 @@ template <typename Group>
 void group_barrier(Group grp, memory_scope fenceScope = Group::fence_scope);
 
 template <>
+inline
 void group_barrier<group<1>>(group<1> grp, memory_scope fenceScope) {
 #ifdef __SYCL_DEVICE_ONLY__
   detail::barrier(
@@ -41,6 +42,7 @@ void group_barrier<group<1>>(group<1> grp, memory_scope fenceScope) {
 }
 
 template <>
+inline
 void group_barrier<group<2>>(group<2> grp, memory_scope fenceScope) {
 #ifdef __SYCL_DEVICE_ONLY__
   detail::barrier(
@@ -51,6 +53,7 @@ void group_barrier<group<2>>(group<2> grp, memory_scope fenceScope) {
 }
 
 template <>
+inline
 void group_barrier<group<3>>(group<3> grp, memory_scope fenceScope) {
 #ifdef __SYCL_DEVICE_ONLY__
   detail::barrier(
@@ -61,6 +64,7 @@ void group_barrier<group<3>>(group<3> grp, memory_scope fenceScope) {
 }
 
 template <>
+inline
 void group_barrier<experimental::sub_group>(experimental::sub_group subGrp,
                                             memory_scope fenceScope) {
 #ifdef __SYCL_DEVICE_ONLY__

This lets the compilation and linking complete properly.

Thanks for the report and solution Andrea, we are looking into this.

Hi Andrea, we’ve patched a fix for this problem and all going to plan it will be in the next release. I will confirm nearer the time.

Hi Rod,
that’s great to hear, thanks !

Ciao,
.Andrea