[ESIMD] Get rid of __SYCL_EXPLICIT_SIMD__ macro (#3404)

Author: DenisBakhvalov

clang/lib/Frontend/InitPreprocessor.cpp

@@ -1139,8 +1139,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
         Builder.defineMacro("__SYCL_NVPTX__", "1");
     }
   }
-  if (LangOpts.SYCLExplicitSIMD)
-    Builder.defineMacro("__SYCL_EXPLICIT_SIMD__", "1");
   if (LangOpts.SYCLUnnamedLambda)
     Builder.defineMacro("__SYCL_UNNAMED_LAMBDA__", "1");
 

sycl/include/CL/sycl/accessor.hpp

@@ -461,9 +461,9 @@ class image_accessor
 private:
   friend class sycl::INTEL::gpu::AccessorPrivateProxy;
 
-#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
+#ifdef __SYCL_DEVICE_ONLY__
   const OCLImageTy getNativeImageObj() const { return MImageObj; }
-#endif // __SYCL_DEVICE_ONLY__ && __SYCL_EXPLICIT_SIMD__
+#endif // __SYCL_DEVICE_ONLY__
 
 public:
   using value_type = DataT;
@@ -858,19 +858,8 @@ class accessor :
 
   detail::AccessorImplDevice<AdjustedDim> impl;
 
-#ifdef __SYCL_EXPLICIT_SIMD__
-  // TODO all the Image1dBuffer* stuff, including the union with MData field
-  // below is not used anymore and is left temporarily to avoid ABI breaking
-  // changes.
-  using OCLImage1dBufferTy =
-      typename detail::opencl_image1d_buffer_type<AccessMode>::type;
-#endif // __SYCL_EXPLICIT_SIMD__
-
   union {
     ConcreteASPtrType MData;
-#ifdef __SYCL_EXPLICIT_SIMD__
-    OCLImage1dBufferTy ImageBuffer;
-#endif // __SYCL_EXPLICIT_SIMD__
   };
 
   // TODO replace usages with getQualifiedPtr

sycl/include/CL/sycl/detail/image_ocl_types.hpp

@@ -192,33 +192,6 @@ inline int getSPIRVElementSize(int ImageChannelType, int ImageChannelOrder) {
   }
 }
 
-#ifdef __SYCL_EXPLICIT_SIMD__
-// TODO all the opencl_image1d_buffer* stuff below is not used anymore and is
-// left temporarily to avoid ABI breaking changes - field of this type is
-// temporarily present in the accessor class.
-template <access::mode AccessMode> struct opencl_image1d_buffer_type;
-
-// OpenCL types used only when compiling DPCPP ESIMD kernels
-#define __SYCL_IMAGE_BUFFER_TY_DEFINE(AccessMode, AMSuffix)                    \
-  template <> struct opencl_image1d_buffer_type<access::mode::AccessMode> {    \
-    using type = __ocl_image1d_buffer_##AMSuffix##_t;                          \
-  }
-
-__SYCL_IMAGE_BUFFER_TY_DEFINE(read, ro);
-__SYCL_IMAGE_BUFFER_TY_DEFINE(write, wo);
-__SYCL_IMAGE_BUFFER_TY_DEFINE(discard_write, wo);
-__SYCL_IMAGE_BUFFER_TY_DEFINE(read_write, rw);
-
-template <> struct opencl_image1d_buffer_type<access::mode::atomic> {
-  // static_assert(false && "atomic access not supported for image1d
-  // buffers");
-  // TODO this should be disabled; currently there is instantiation of this
-  // class happenning even if atomic access not used - using dummy type
-  // definition for now.
-  using type = unsigned int;
-};
-#endif // __SYCL_EXPLICIT_SIMD__
-
 template <int Dimensions, access::mode AccessMode, access::target AccessTarget>
 struct opencl_image_type;