[SYCL] Remove the INLINE_IF_DEVICE workaround

sergey-semenov · vladimirlaz · commit e8891d7eb576 · 2019-02-12T07:44:06.000+03:00
Compilation flow has switched from linking SPIRV to linking
LLVM bytecode. This workaround is no longer needed.

Signed-off-by: Semenov, Sergey &lt;sergey.semenov@intel.com&gt;
Signed-off-by: Vladimir Lazarev &lt;vladimir.lazarev@intel.com&gt;
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
@@ -43,7 +43,7 @@ class subscript_obj {
   subscript_obj(const accessor_t &acc, cl::sycl::id<accessorDim> &indexes)
       : accRef(acc), ids(indexes) {}
 
-  INLINE_IF_DEVICE subscript_obj<accessorDim, dataT, dimensions - 1, accessMode, accessTarget,
+  subscript_obj<accessorDim, dataT, dimensions - 1, accessMode, accessTarget,
                 isPlaceholder>
   operator[](size_t index) {
     ids[accessorDim - dimensions] = index;
@@ -66,7 +66,7 @@ class subscript_obj<accessorDim, dataT, 1, accessMode, accessTarget,
   subscript_obj(const accessor_t &acc, cl::sycl::id<accessorDim> &indexes)
       : accRef(acc), ids(indexes) {}
 
-  INLINE_IF_DEVICE dataT &operator[](size_t index) {
+  dataT &operator[](size_t index) {
     ids[accessorDim - 1] = index;
     return accRef.__impl()->Data[getOffsetForId(
       accRef.__impl()->Range, ids, accRef.__impl()->Offset)];
@@ -87,7 +87,7 @@ class subscript_obj<accessorDim, dataT, 1, access::mode::read, accessTarget,
   subscript_obj(const accessor_t &acc, cl::sycl::id<accessorDim> &indexes)
       : accRef(acc), ids(indexes) {}
 
-  INLINE_IF_DEVICE typename detail::remove_AS<dataT>::type
+  typename detail::remove_AS<dataT>::type
   operator[](size_t index) {
     ids[accessorDim - 1] = index;
     return accRef.__impl()->Data[getOffsetForId(
@@ -118,7 +118,7 @@ SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions == 0) {
   accessor_impl(dataT *Data) : Data(Data) {}
 
   // Returns the number of accessed elements.
-  INLINE_IF_DEVICE size_t get_count() const { return 1; }
+  size_t get_count() const { return 1; }
 };
 
 /// Implementation of host accessor.
@@ -133,7 +133,7 @@ SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions > 0) {
       : Data(Data), Range(Range), Offset(Offset) {}
 
   // Returns the number of accessed elements.
-  INLINE_IF_DEVICE size_t get_count() const { return Range.size(); }
+  size_t get_count() const { return Range.size(); }
 };
 
 /// Implementation of device (kernel) accessor providing access to a single
@@ -163,7 +163,7 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
   {}
 
   // Returns the number of accessed elements.
-  INLINE_IF_DEVICE size_t get_count() const { return 1; }
+  size_t get_count() const { return 1; }
 
   static_assert(
       std::is_same<typename DeviceValueType<dataT, accessTarget>::type,
@@ -201,7 +201,7 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
   {}
 
   // Returns the number of accessed elements.
-  INLINE_IF_DEVICE size_t get_count() const { return Range.size(); }
+  size_t get_count() const { return Range.size(); }
 
   static_assert(
       std::is_same<typename DeviceValueType<dataT, accessTarget>::type,
@@ -240,7 +240,7 @@ SYCL_ACCESSOR_IMPL(accessTarget == access::target::local &&
   }
 
   // Returns the number of accessed elements.
-  INLINE_IF_DEVICE size_t get_count() const { return 1; }
+  size_t get_count() const { return 1; }
 
   static_assert(
       std::is_same<typename DeviceValueType<dataT, accessTarget>::type,
@@ -285,7 +285,7 @@ SYCL_ACCESSOR_IMPL(accessTarget == access::target::local &&
   }
 
   // Returns the number of accessed elements.
-  INLINE_IF_DEVICE size_t get_count() const { return Range.size(); }
+  size_t get_count() const { return Range.size(); }
 
   static_assert(
       std::is_same<typename DeviceValueType<dataT, accessTarget>::type,
@@ -307,11 +307,11 @@ class accessor_base {
   using _ImplT =
       accessor_impl<dataT, dimensions, accessMode, accessTarget, isPlaceholder>;
 
-  INLINE_IF_DEVICE const _ImplT *__impl() const {
+  const _ImplT *__impl() const {
     return reinterpret_cast<const _ImplT *>(this);
   }
 
-  INLINE_IF_DEVICE _ImplT *__impl() { return reinterpret_cast<_ImplT *>(this); }
+  _ImplT *__impl() { return reinterpret_cast<_ImplT *>(this); }
 
   static_assert(
       std::is_same<typename DeviceValueType<dataT, accessTarget>::type,
@@ -339,21 +339,21 @@ class accessor_base {
 
 SYCL_ACCESSOR_SUBCLASS(accessor_common, accessor_base, true /* always */) {
   // Returns true if the current accessor is a placeholder accessor.
-  INLINE_IF_DEVICE constexpr bool is_placeholder() const {
+  constexpr bool is_placeholder() const {
     return isPlaceholder == access::placeholder::true_t;
   }
 
   // Returns the size of the accessed memory in bytes.
-  INLINE_IF_DEVICE size_t get_size() const { return this->get_count() * sizeof(dataT); }
+  size_t get_size() const { return this->get_count() * sizeof(dataT); }
 
   // Returns the number of accessed elements.
-  INLINE_IF_DEVICE size_t get_count() const { return this->__impl()->get_count(); }
+  size_t get_count() const { return this->__impl()->get_count(); }
 
-  template <int Dimensions = dimensions> INLINE_IF_DEVICE
+  template <int Dimensions = dimensions>
   typename std::enable_if<(Dimensions > 0), range<Dimensions>>::type
   get_range() const { return this->__impl()->Range; }
 
-  template <int Dimensions = dimensions> INLINE_IF_DEVICE
+  template <int Dimensions = dimensions>
   typename std::enable_if<(Dimensions > 0), id<Dimensions>>::type
   get_offset() const { return this->__impl()->Offset; }
 };
@@ -364,7 +364,7 @@ SYCL_ACCESSOR_SUBCLASS(accessor_opdata_w, accessor_common,
                         accessMode == access::mode::discard_write ||
                         accessMode == access::mode::discard_read_write) &&
                        dimensions == 0) {
-  INLINE_IF_DEVICE operator dataT &() const {
+  operator dataT &() const {
     return this->__impl()->Data[0];
   }
 };
@@ -382,7 +382,7 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_wn, accessor_opdata_w,
 
   subscript_obj<dimensions, dataT, dimensions - 1, accessMode, accessTarget,
               isPlaceholder>
-  INLINE_IF_DEVICE operator[](size_t index) const {
+  operator[](size_t index) const {
     id<dimensions> ids;
     ids[0] = index;
     return subscript_obj<dimensions, dataT, dimensions - 1, accessMode,
@@ -402,11 +402,11 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_w, accessor_subscript_wn,
   // classes. That's why operator[] defined in accessor_subscript_wn
   // is not visible here and we have to define
   // operator[](id<dimensions>) once again.
-  INLINE_IF_DEVICE dataT &operator[](id<dimensions> index) const {
+  dataT &operator[](id<dimensions> index) const {
     return this->operator[](
       getOffsetForId(this->get_range(), index, this->get_offset()));
   }
-  INLINE_IF_DEVICE dataT &operator[](size_t index) const {
+  dataT &operator[](size_t index) const {
     return this->__impl()->Data[index];
   }
 };
@@ -560,7 +560,7 @@ class accessor
   // implementation.
   _ImplT __impl;
 
-  INLINE_IF_DEVICE void __init(_ValueType *Ptr, range<dimensions> Range,
+  void __init(_ValueType *Ptr, range<dimensions> Range,
       id<dimensions> Offset) {
     __impl.Data = Ptr;
     __impl.Range = Range;
diff --git a/sycl/include/CL/sycl/detail/array.hpp b/sycl/include/CL/sycl/detail/array.hpp
@@ -21,67 +21,67 @@ namespace detail {
 
 template <int dimensions = 1> class array {
 public:
-  INLINE_IF_DEVICE array() : common_array{0} {}
+  array() : common_array{0} {}
 
   /* The following constructor is only available in the array struct
    * specialization where: dimensions==1 */
-  template <int N = dimensions> INLINE_IF_DEVICE
+  template <int N = dimensions>
   array(typename std::enable_if<(N == 1), size_t>::type dim0)
       : common_array{dim0} {}
 
   /* The following constructor is only available in the array struct
    * specialization where: dimensions==2 */
-  template <int N = dimensions> INLINE_IF_DEVICE
+  template <int N = dimensions>
   array(typename std::enable_if<(N == 2), size_t>::type dim0, size_t dim1)
       : common_array{dim0, dim1} {}
 
   /* The following constructor is only available in the array struct
    * specialization where: dimensions==3 */
-  template <int N = dimensions> INLINE_IF_DEVICE
+  template <int N = dimensions>
   array(typename std::enable_if<(N == 3), size_t>::type dim0, size_t dim1,
         size_t dim2)
       : common_array{dim0, dim1, dim2} {}
 
   // Conversion operators to derived classes
-  INLINE_IF_DEVICE operator cl::sycl::id<dimensions>() const {
+  operator cl::sycl::id<dimensions>() const {
     cl::sycl::id<dimensions> result;
     for (int i = 0; i < dimensions; ++i) {
       result[i] = common_array[i];
     }
     return result;
   }
 
-  INLINE_IF_DEVICE operator cl::sycl::range<dimensions>() const {
+  operator cl::sycl::range<dimensions>() const {
     cl::sycl::range<dimensions> result;
     for (int i = 0; i < dimensions; ++i) {
       result[i] = common_array[i];
     }
     return result;
   }
 
-  INLINE_IF_DEVICE size_t get(int dimension) const {
+  size_t get(int dimension) const {
     check_dimension(dimension);
     return common_array[dimension];
   }
 
-  INLINE_IF_DEVICE size_t &operator[](int dimension) {
+  size_t &operator[](int dimension) {
     check_dimension(dimension);
     return common_array[dimension];
   }
 
-  INLINE_IF_DEVICE size_t operator[](int dimension) const {
+  size_t operator[](int dimension) const {
     check_dimension(dimension);
     return common_array[dimension];
   }
 
-  INLINE_IF_DEVICE array(const array<dimensions> &rhs) = default;
-  INLINE_IF_DEVICE array(array<dimensions> &&rhs) = default;
-  INLINE_IF_DEVICE array<dimensions> &operator=(const array<dimensions> &rhs) = default;
-  INLINE_IF_DEVICE array<dimensions> &operator=(array<dimensions> &&rhs) = default;
+  array(const array<dimensions> &rhs) = default;
+  array(array<dimensions> &&rhs) = default;
+  array<dimensions> &operator=(const array<dimensions> &rhs) = default;
+  array<dimensions> &operator=(array<dimensions> &&rhs) = default;
 
   // Returns true iff all elements in 'this' are equal to
   // the corresponding elements in 'rhs'.
-  INLINE_IF_DEVICE bool operator==(const array<dimensions> &rhs) const {
+  bool operator==(const array<dimensions> &rhs) const {
     for (int i = 0; i < dimensions; ++i) {
       if (this->common_array[i] != rhs.common_array[i]) {
         return false;
@@ -92,7 +92,7 @@ template <int dimensions = 1> class array {
 
   // Returns true iff there is at least one element in 'this'
   // which is not equal to the corresponding element in 'rhs'.
-  INLINE_IF_DEVICE bool operator!=(const array<dimensions> &rhs) const {
+  bool operator!=(const array<dimensions> &rhs) const {
     for (int i = 0; i < dimensions; ++i) {
       if (this->common_array[i] != rhs.common_array[i]) {
         return true;
diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp
@@ -74,16 +74,6 @@ const char *stringifyErrorCode(cl_int error);
 #define ALWAYS_INLINE
 #endif
 
-// TODO this macro is introduced to workaround SPIRV translator problem with
-// dropping linkonce_odr attribute leading to duplicated symbol errors in
-// the bitcode linker for functions defined in the headers. Remove once fixed.
-#ifdef __SYCL_DEVICE_ONLY__
-#define INLINE_IF_DEVICE ALWAYS_INLINE
-#else
-#define INLINE_IF_DEVICE
-#endif // __SYCL_DEVICE_ONLY__
-
-
 namespace cl {
 namespace sycl {
 namespace detail {
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
@@ -61,20 +61,20 @@ namespace detail {
   template <int Dim, class DstT> struct InitSizesST##POSTFIX;                  \
                                                                                \
   template <class DstT> struct InitSizesST##POSTFIX<1, DstT> {                 \
-    static INLINE_IF_DEVICE void initSize(DstT &Dst) {                         \
+    static void initSize(DstT &Dst) {                                          \
       Dst[0] = cl::__spirv::get##POSTFIX<0>();                                 \
     }                                                                          \
   };                                                                           \
                                                                                \
   template <class DstT> struct InitSizesST##POSTFIX<2, DstT> {                 \
-    static INLINE_IF_DEVICE void initSize(DstT &Dst) {                         \
+    static void initSize(DstT &Dst) {                                          \
       Dst[1] = cl::__spirv::get##POSTFIX<1>();                                 \
       InitSizesST##POSTFIX<1, DstT>::initSize(Dst);                            \
     }                                                                          \
   };                                                                           \
                                                                                \
   template <class DstT> struct InitSizesST##POSTFIX<3, DstT> {                 \
-    static INLINE_IF_DEVICE void initSize(DstT &Dst) {                         \
+    static void initSize(DstT &Dst) {                                          \
       Dst[2] = cl::__spirv::get##POSTFIX<2>();                                 \
       InitSizesST##POSTFIX<2, DstT>::initSize(Dst);                            \
     }                                                                          \
diff --git a/sycl/include/CL/sycl/id.hpp b/sycl/include/CL/sycl/id.hpp
@@ -19,7 +19,7 @@ template <int dimensions> class range;
 template <int dimensions = 1> struct id : public detail::array<dimensions> {
 public:
   using base = detail::array<dimensions>;
-  INLINE_IF_DEVICE id() = default;
+  id() = default;
 
   /* The following constructor is only available in the id struct
    * specialization where: dimensions==1 */
@@ -149,7 +149,7 @@ template <int dimensions = 1> struct id : public detail::array<dimensions> {
 };
 
 namespace detail {
-template <int dimensions> INLINE_IF_DEVICE
+template <int dimensions>
 size_t getOffsetForId(range<dimensions> Range, id<dimensions> Id,
                       id<dimensions> Offset) {
   size_t offset = 0;
diff --git a/sycl/test/separate-compile/test.cpp b/sycl/test/separate-compile/test.cpp
@@ -1,30 +1,25 @@
 // >> ---- compile src1
 // >> device compilation...
-// RUN: %clang -std=c++11 -fno-sycl-use-bitcode --sycl -Xclang -fsycl-int-header=sycl_ihdr_a.h %s -c -o a_kernel.spv
+// RUN: %clang -std=c++11 --sycl -Xclang -fsycl-int-header=sycl_ihdr_a.h %s -c -o a_kernel.bc
 // >> host compilation...
 // RUN: %clang -std=c++11 -include sycl_ihdr_a.h -g -c %s -o a.o
 //
 // >> ---- compile src2
 // >> device compilation...
-// RUN: %clang -DB_CPP=1 -std=c++11 -fno-sycl-use-bitcode --sycl -Xclang -fsycl-int-header=sycl_ihdr_b.h %s -c -o b_kernel.spv
+// RUN: %clang -DB_CPP=1 -std=c++11 --sycl -Xclang -fsycl-int-header=sycl_ihdr_b.h %s -c -o b_kernel.bc
 // >> host compilation...
 // RUN: %clang -DB_CPP=1 -std=c++11 -include sycl_ihdr_b.h -g -c %s -o b.o
 //
 // >> ---- bundle .o with .spv
 // >> run bundler
-// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=a.o,a_kernel.spv -outputs=a_fat.o
-// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=b.o,b_kernel.spv -outputs=b_fat.o
+// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=a.o,a_kernel.bc -outputs=a_fat.o
+// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=b.o,b_kernel.bc -outputs=b_fat.o
 //
 // >> ---- unbundle fat objects
-// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=a.o,a_kernel.spv -inputs=a_fat.o -unbundle
-// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=b.o,b_kernel.spv -inputs=b_fat.o -unbundle
+// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=a.o,a_kernel.bc -inputs=a_fat.o -unbundle
+// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=b.o,b_kernel.bc -inputs=b_fat.o -unbundle
 //
 // >> ---- link device code
-// >> convert to bitcode
-// RUN: llvm-spirv -r -o=a_kernel.bc a_kernel.spv
-// RUN: llvm-spirv -r -o=b_kernel.bc b_kernel.spv
-//
-// >> link bitcode
 // RUN: llvm-link -o=app.bc a_kernel.bc b_kernel.bc
 //
 // >> convert linked .bc to spirv
