Inlining atomic_fence instead of template argument for odr failure

Author: dongkyunahn-intel

sycl/include/sycl/ext/intel/esimd/detail/atomic_intrin.hpp

@@ -205,8 +205,7 @@ template <typename Ty> Ty atomic_cmpxchg(Ty *ptr, Ty desired, Ty expected) {
 #endif
 }
 
-// Template argument is to avoid failure from 'odr.cpp' test
-template <int N> void atomic_fence() {
+inline void atomic_fence() {
 #ifdef _WIN32
   // TODO: Windows will be supported soon
   __ESIMD_UNSUPPORTED_ON_HOST;

sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp

@@ -616,7 +616,7 @@ __ESIMD_INTRIN void __esimd_fence(uint8_t cntl)
 {
   // CM_EMU's 'cm_fence' is NOP. Disabled.
   // sycl::detail::getESIMDDeviceInterface()->cm_fence_ptr();
-  __ESIMD_DNS::atomic_fence<0>();
+  __ESIMD_DNS::atomic_fence();
 }
 #endif // __SYCL_DEVICE_ONLY__
 

sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp

@@ -1808,7 +1808,7 @@ __ESIMD_INTRIN void __esimd_lsc_fence(__ESIMD_DNS::simd_mask_storage_t<N> pred)
     ;
 #else  // __SYCL_DEVICE_ONLY__
 {
-  __ESIMD_DNS::atomic_fence<N>();
+  __ESIMD_DNS::atomic_fence();
 }
 #endif // __SYCL_DEVICE_ONLY__