Skip to content

[ESIMD] Get rid of __SYCL_EXPLICIT_SIMD__ macro #3404

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 6 commits into from
Mar 30, 2021
Merged

[ESIMD] Get rid of __SYCL_EXPLICIT_SIMD__ macro #3404

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
2f09dbd
[ESIMD] Get rid of __SYCL_EXPLICIT_SIMD__ macro
DenisBakhvalov Mar 23, 2021
eb0d516
Fixed failing test
DenisBakhvalov Mar 24, 2021
68bd321
Fixed failing llvm-test-suite accessor test
DenisBakhvalov Mar 24, 2021
7910a47
Revert "Fixed failing test"
DenisBakhvalov Mar 25, 2021
d0d57c1
Break ABI of ESIMD instead of SYCL device code
DenisBakhvalov Mar 25, 2021
0ff2190
Cleaned up other ESIMD leftover code
DenisBakhvalov Mar 25, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 0 additions & 2 deletions clang/lib/Frontend/InitPreprocessor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1133,8 +1133,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");

Expand Down
15 changes: 2 additions & 13 deletions sycl/include/CL/sycl/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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
Expand Down
27 changes: 0 additions & 27 deletions sycl/include/CL/sycl/detail/image_ocl_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand Down