diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 7e9fbd6105b78..8a8136418c7a3 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -15,7 +15,7 @@ namespace sycl { namespace access { enum class target { - global_buffer = 2014, + global_buffer __SYCL2020_DEPRECATED("use 'target::device' instead") = 2014, constant_buffer = 2015, local = 2016, image = 2017, @@ -137,7 +137,7 @@ template struct TargetToAS { }; #ifdef __ENABLE_USM_ADDR_SPACE__ -template <> struct TargetToAS { +template <> struct TargetToAS { constexpr static access::address_space AS = access::address_space::global_device_space; }; diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index b5048c60538e6..5430b465a849d 100755 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -219,7 +219,7 @@ class AccessorPrivateProxy; template > class accessor; @@ -271,6 +271,9 @@ class accessor_common { constexpr static bool IsHostBuf = AccessTarget == access::target::host_buffer; + // TODO: SYCL 2020 deprecates four of the target enum values + // and replaces them with 2 (device and host_task). May want + // to change these constexpr. constexpr static bool IsGlobalBuf = AccessTarget == access::target::global_buffer; @@ -1605,9 +1608,9 @@ class accessor : return getQualifiedPtr() + LinearIndex; } - template > + template < + access::target AccessTarget_ = AccessTarget, + typename = detail::enable_if_t> global_ptr get_pointer() const { const size_t LinearIndex = getLinearIndex(id()); return global_ptr(getQualifiedPtr() + LinearIndex); @@ -1638,21 +1641,21 @@ class accessor : template accessor(buffer) - -> accessor; + -> accessor; template accessor(buffer, const ext::oneapi::accessor_property_list &) - -> accessor accessor>; template accessor(buffer, Type1) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t>; template , Type1, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t, ext::oneapi::accessor_property_list>; @@ -1668,7 +1671,7 @@ template accessor(buffer, Type1, Type2) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t>; template , Type1, Type2, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t, ext::oneapi::accessor_property_list>; @@ -1684,7 +1687,7 @@ template accessor(buffer, Type1, Type2, Type3) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t>; template , Type1, Type2, Type3, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t, ext::oneapi::accessor_property_list>; @@ -1700,7 +1703,7 @@ template accessor(buffer, Type1, Type2, Type3, Type4) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t>; template , Type1, Type2, Type3, Type4, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::true_t, ext::oneapi::accessor_property_list>; template accessor(buffer, handler) - -> accessor; + -> accessor; template accessor(buffer, handler, const ext::oneapi::accessor_property_list &) - -> accessor accessor>; template accessor(buffer, handler, Type1) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t>; template , handler, Type1, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t, ext::oneapi::accessor_property_list>; @@ -1744,7 +1747,7 @@ template accessor(buffer, handler, Type1, Type2) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t>; template , handler, Type1, Type2, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t, ext::oneapi::accessor_property_list>; @@ -1760,7 +1763,7 @@ template accessor(buffer, handler, Type1, Type2, Type3) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t>; template , handler, Type1, Type2, Type3, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t, ext::oneapi::accessor_property_list>; @@ -1777,7 +1780,7 @@ template , handler, Type1, Type2, Type3, Type4) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t>; template , handler, Type1, Type2, Type3, Type4, const ext::oneapi::accessor_property_list &) -> accessor(), - detail::deduceAccessTarget(target::global_buffer), + detail::deduceAccessTarget(target::device), access::placeholder::false_t, ext::oneapi::accessor_property_list>; #endif diff --git a/sycl/include/CL/sycl/backend/cuda.hpp b/sycl/include/CL/sycl/backend/cuda.hpp index 9a45c2fc8bf0e..08a2b80c8a073 100644 --- a/sycl/include/CL/sycl/backend/cuda.hpp +++ b/sycl/include/CL/sycl/backend/cuda.hpp @@ -44,9 +44,9 @@ template <> struct interop { using type = CUevent; }; template <> struct interop { using type = CUmodule; }; template -struct interop> { +struct interop> { using type = CUdeviceptr; }; diff --git a/sycl/include/CL/sycl/backend/opencl.hpp b/sycl/include/CL/sycl/backend/opencl.hpp index c051ca0f0a4ce..6854484cd6e08 100644 --- a/sycl/include/CL/sycl/backend/opencl.hpp +++ b/sycl/include/CL/sycl/backend/opencl.hpp @@ -42,9 +42,9 @@ template <> struct interop { template <> struct interop { using type = cl_event; }; template -struct interop> { +struct interop> { using type = cl_mem; }; diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp index 240c332c910b4..eee3df0d94eac 100755 --- a/sycl/include/CL/sycl/buffer.hpp +++ b/sycl/include/CL/sycl/buffer.hpp @@ -272,8 +272,7 @@ class buffer { return impl->template get_allocator(); } - template + template accessor> get_access(handler &CommandGroupHandler) { @@ -291,8 +290,7 @@ class buffer { ext::oneapi::accessor_property_list<>>(*this); } - template + template accessor> get_access(handler &commandGroupHandler, range accessRange, diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index a6e4af3da5336..2f0c989c07b78 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -751,7 +751,7 @@ class __SYCL_EXPORT handler { #endif // __SYCL_DEVICE_ONLY__ constexpr static bool isConstOrGlobal(access::target AccessTarget) { - return AccessTarget == access::target::global_buffer || + return AccessTarget == access::target::device || AccessTarget == access::target::constant_buffer; } diff --git a/sycl/include/CL/sycl/interop_handle.hpp b/sycl/include/CL/sycl/interop_handle.hpp index 27d3f6195f638..54b6abfd3b9f2 100644 --- a/sycl/include/CL/sycl/interop_handle.hpp +++ b/sycl/include/CL/sycl/interop_handle.hpp @@ -52,9 +52,7 @@ class interop_handle { access::mode Mode, access::target Target, access::placeholder IsPlh> backend_return_t> get_native_mem(const accessor &Acc) const { - // TODO: the method is available when the target is target::device. Add it - // to the assert below when target::device enum is created. - static_assert(Target == access::target::global_buffer || + static_assert(Target == access::target::device || Target == access::target::constant_buffer, "The method is available only for target::device accessors"); #ifndef __SYCL_DEVICE_ONLY__ diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index d9b19a4db1c4d..07c506897b930 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -117,10 +117,9 @@ template class multi_ptr { _Space == Space && (Space == access::address_space::global_space || Space == access::address_space::global_device_space)>> - multi_ptr( - accessor - Accessor) { + multi_ptr(accessor + Accessor) { m_Pointer = (pointer_t)(Accessor.get_pointer().get()); } @@ -166,10 +165,9 @@ template class multi_ptr { (Space == access::address_space::global_space || Space == access::address_space::global_device_space) && std::is_const::value && std::is_same::value>> - multi_ptr( - accessor, dimensions, Mode, - access::target::global_buffer, isPlaceholder, PropertyListT> - Accessor) + multi_ptr(accessor, dimensions, Mode, + access::target::device, isPlaceholder, PropertyListT> + Accessor) : multi_ptr(Accessor.get_pointer()) {} // Only if Space == local_space and element type is const @@ -382,10 +380,9 @@ template class multi_ptr { _Space == Space && (Space == access::address_space::global_space || Space == access::address_space::global_device_space)>> - multi_ptr( - accessor - Accessor) + multi_ptr(accessor + Accessor) : multi_ptr(Accessor.get_pointer()) {} // Only if Space == local_space @@ -503,10 +500,9 @@ class multi_ptr { _Space == Space && (Space == access::address_space::global_space || Space == access::address_space::global_device_space)>> - multi_ptr( - accessor - Accessor) + multi_ptr(accessor + Accessor) : multi_ptr(Accessor.get_pointer()) {} // Only if Space == local_space @@ -557,9 +553,9 @@ class multi_ptr { #ifdef __cpp_deduction_guides template -multi_ptr(accessor) - ->multi_ptr; +multi_ptr(accessor) + -> multi_ptr; template multi_ptr(accessor system_scope{}; template class atomic_accessor : public accessor atomic_accessor(buffer, order_tag_t, scope_tag_t, property_list = {}) - -> atomic_accessor atomic_accessor; template atomic_accessor(buffer, handler, order_tag_t, scope_tag_t, property_list = {}) - -> atomic_accessor atomic_accessor; #endif diff --git a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp index 412f944a4ac0c..2301766b51cfe 100644 --- a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp +++ b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp @@ -40,9 +40,9 @@ template <> struct interop { }; template -struct interop> { +struct interop> { using type = char *; }; diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp index 7d3974f593f80..b80d3d9bf0cdf 100644 --- a/sycl/include/sycl/ext/oneapi/reduction.hpp +++ b/sycl/include/sycl/ext/oneapi/reduction.hpp @@ -376,12 +376,11 @@ class reduction_impl : private reduction_impl_base { using result_type = T; using binary_operation = BinaryOperation; using rw_accessor_type = - accessor>; using dw_accessor_type = - accessor>; + accessor>; static constexpr int accessor_dim = Dims; static constexpr int buffer_dim = (Dims == 0) ? 1 : Dims; @@ -625,9 +624,9 @@ class reduction_impl : private reduction_impl_base { void associateWithHandler(handler &CGH) { #ifndef __SYCL_DEVICE_ONLY__ if (MRWAcc) - CGH.associateWithHandler(MRWAcc.get(), access::target::global_buffer); + CGH.associateWithHandler(MRWAcc.get(), access::target::device); else if (MDWAcc) - CGH.associateWithHandler(MDWAcc.get(), access::target::global_buffer); + CGH.associateWithHandler(MDWAcc.get(), access::target::device); #else (void)CGH; #endif @@ -714,7 +713,7 @@ class reduction_impl : private reduction_impl_base { return createHandlerWiredReadWriteAccessor(CGH, *MOutBufPtr); } - accessor getReadWriteAccessorToInitializedGroupsCounter(handler &CGH) { auto CounterMem = std::make_shared(0); @@ -2009,7 +2008,7 @@ tuple_select_elements(TupleT Tuple, std::index_sequence) { template detail::reduction_impl -reduction(accessor &Acc, +reduction(accessor &Acc, const T &Identity, BinaryOperation BOp) { return {Acc, Identity, BOp}; } @@ -2022,7 +2021,7 @@ template std::enable_if_t::value, detail::reduction_impl> -reduction(accessor &Acc, +reduction(accessor &Acc, BinaryOperation) { return {Acc}; } diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index abbf84e17aaa5..7f2b0f14dc7da 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -341,7 +341,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, // The first 11 bits of Size encodes the accessor target. const access::target AccTarget = static_cast(Size & 0x7ff); switch (AccTarget) { - case access::target::global_buffer: + case access::target::device: case access::target::constant_buffer: { detail::Requirement *AccImpl = static_cast(Ptr); addArgsForGlobalAccessor(AccImpl, Index, IndexShift, Size, @@ -468,7 +468,7 @@ void handler::extractArgsAndReqsFromLambda( // The first 11 bits of Size encodes the accessor target. const access::target AccTarget = static_cast(Size & 0x7ff); - if ((AccTarget == access::target::global_buffer || + if ((AccTarget == access::target::device || AccTarget == access::target::constant_buffer) || (AccTarget == access::target::image || AccTarget == access::target::image_array)) { diff --git a/sycl/test/abi/symbol_size_alignment.cpp b/sycl/test/abi/symbol_size_alignment.cpp index 4ff2cdac83d8e..94582e9e595aa 100644 --- a/sycl/test/abi/symbol_size_alignment.cpp +++ b/sycl/test/abi/symbol_size_alignment.cpp @@ -40,9 +40,9 @@ template void check() { } int main() { - using accessor_t = accessor; + using accessor_t = + accessor; check(); check, 24, 8>(); check, 24, 8>(); diff --git a/sycl/test/abi/user_mangling.cpp b/sycl/test/abi/user_mangling.cpp index 8b1c83e3953cd..3d41f5d8f9bd7 100644 --- a/sycl/test/abi/user_mangling.cpp +++ b/sycl/test/abi/user_mangling.cpp @@ -7,6 +7,9 @@ #include #ifdef __SYCL_DEVICE_ONLY__ +// CHK-DEVICE: define dso_local spir_func void @_Z4accdN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}}) +SYCL_EXTERNAL void accd(sycl::accessor) {} + // CHK-DEVICE: define dso_local spir_func void @_Z3accN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}}) SYCL_EXTERNAL void acc(sycl::accessor) {} diff --git a/sycl/test/basic_tests/accessor/accessor_property_list_ct.cpp b/sycl/test/basic_tests/accessor/accessor_property_list_ct.cpp index e4d791d5c16dc..dc87884db0588 100644 --- a/sycl/test/basic_tests/accessor/accessor_property_list_ct.cpp +++ b/sycl/test/basic_tests/accessor/accessor_property_list_ct.cpp @@ -7,7 +7,7 @@ using namespace sycl::ext::oneapi; void foo(sycl::accessor, property::no_offset::instance<>>> @@ -111,7 +111,7 @@ int main() { accessor_property_list PL{sycl::no_init, no_alias}; sycl::accessor acc_1(buf_data, PL); sycl::accessor>> acc_2(acc_1); diff --git a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp index ebc2540ee44d9..0e793032f60e5 100644 --- a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp +++ b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp @@ -22,6 +22,9 @@ int main() { buffer ConstantBuf(Range); queue q; q.submit([&](handler &Cgh) { + auto DeviceRWAcc = + GlobalBuf.get_access(Cgh); + auto DeviceRAcc = GlobalBuf.get_access(Cgh); auto GlobalRWAcc = GlobalBuf.get_access(Cgh); auto GlobalRAcc = @@ -31,6 +34,10 @@ int main() { accessor LocalAcc(Range, Cgh); Cgh.single_task([=]() { + static_assert(std::is_same::value, + "Incorrect type from device read-write accessor"); + static_assert(std::is_same::value, + "Incorrect type from device read accessor"); static_assert(std::is_same::value, "Incorrect type from global read-write accessor"); static_assert(std::is_same::value, diff --git a/sycl/test/basic_tests/reduction_ctor.cpp b/sycl/test/basic_tests/reduction_ctor.cpp index 07a4fbd856961..6f673f48b6c66 100644 --- a/sycl/test/basic_tests/reduction_ctor.cpp +++ b/sycl/test/basic_tests/reduction_ctor.cpp @@ -48,11 +48,11 @@ void testKnown(T Identity, BinaryOperation BOp, T A, T B) { T *ReduUSMPtr = malloc_host(1, Q); Q.submit([&](handler &CGH) { - // Reduction needs a global_buffer accessor as a parameter. + // Reduction needs a device accessor as a parameter. // This accessor is not really used in this test. - accessor + accessor ReduRWAcc(ReduBuf, CGH); - accessor + accessor ReduDWAcc(ReduBuf, CGH); auto ReduRW = ext::oneapi::reduction(ReduRWAcc, BOp); auto ReduDW = ext::oneapi::reduction(ReduDWAcc, BOp); @@ -92,11 +92,11 @@ void testUnknown(T Identity, BinaryOperation BOp, T A, T B) { buffer ReduBuf(1); T *ReduUSMPtr = malloc_host(1, Q); Q.submit([&](handler &CGH) { - // Reduction needs a global_buffer accessor as a parameter. + // Reduction needs a device accessor as a parameter. // This accessor is not really used in this test. - accessor + accessor ReduRWAcc(ReduBuf, CGH); - accessor + accessor ReduDWAcc(ReduBuf, CGH); auto ReduRW = ext::oneapi::reduction(ReduRWAcc, Identity, BOp); auto ReduDW = ext::oneapi::reduction(ReduDWAcc, Identity, BOp); diff --git a/sycl/test/check_device_code/buffer_location_codegen.cpp b/sycl/test/check_device_code/buffer_location_codegen.cpp index 61d7ad5c9a6ce..5d23b2498af75 100644 --- a/sycl/test/check_device_code/buffer_location_codegen.cpp +++ b/sycl/test/check_device_code/buffer_location_codegen.cpp @@ -8,8 +8,7 @@ struct Base { int A, B; cl::sycl::accessor< - char, 1, cl::sycl::access::mode::read, - cl::sycl::access::target::global_buffer, + char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::device, cl::sycl::access::placeholder::false_t, cl::sycl::ext::oneapi::accessor_property_list< cl::sycl::ext::intel::property::buffer_location::instance<2>>> @@ -20,7 +19,7 @@ struct Captured : Base, cl::sycl::accessor< char, 1, cl::sycl::access::mode::read, - cl::sycl::access::target::global_buffer, + cl::sycl::access::target::device, cl::sycl::access::placeholder::false_t, cl::sycl::ext::oneapi::accessor_property_list< cl::sycl::ext::intel::property::buffer_location::instance<2>>> { @@ -31,8 +30,7 @@ int main() { Captured Obj; cl::sycl::accessor< int, 1, cl::sycl::access::mode::read_write, - cl::sycl::access::target::global_buffer, - cl::sycl::access::placeholder::false_t, + cl::sycl::access::target::device, cl::sycl::access::placeholder::false_t, cl::sycl::ext::oneapi::accessor_property_list< cl::sycl::ext::intel::property::buffer_location::instance<3>>> accessorA; diff --git a/sycl/test/esimd/block_load_store.cpp b/sycl/test/esimd/block_load_store.cpp index 7b6687e9187b3..8718f539a27e1 100644 --- a/sycl/test/esimd/block_load_store.cpp +++ b/sycl/test/esimd/block_load_store.cpp @@ -10,9 +10,9 @@ using namespace sycl::ext::intel::experimental::esimd; using namespace cl::sycl; -SYCL_EXTERNAL void kernel1( - accessor - &buf) SYCL_ESIMD_FUNCTION { +SYCL_EXTERNAL void +kernel1(accessor &buf) + SYCL_ESIMD_FUNCTION { simd v1(0, 1); auto v0 = block_load(buf, 0); v0 = v0 + v1; diff --git a/sycl/test/esimd/flat_atomic.cpp b/sycl/test/esimd/flat_atomic.cpp index d9eeb4f7265af..bc9d3fde2fb0e 100644 --- a/sycl/test/esimd/flat_atomic.cpp +++ b/sycl/test/esimd/flat_atomic.cpp @@ -12,7 +12,7 @@ using namespace sycl::ext::intel::experimental::esimd; using namespace cl::sycl; void kernel0(accessor &buf) SYCL_ESIMD_FUNCTION { + access::target::device> &buf) SYCL_ESIMD_FUNCTION { simd offsets(0, 1); // CHECK: flat_atomic.cpp:20{{.*}}warning: 'ATOMIC_INC' is deprecated @@ -23,7 +23,7 @@ void kernel0(accessor &buf) SYCL_ESIMD_FUNCTION { + access::target::device> &buf) SYCL_ESIMD_FUNCTION { simd offsets(0, 1); simd v1(0, 1); @@ -35,7 +35,7 @@ void kernel1(accessor &buf) SYCL_ESIMD_FUNCTION { + access::target::device> &buf) SYCL_ESIMD_FUNCTION { simd offsets(0, 1); simd v1(0, 1); diff --git a/sycl/test/esimd/gather_scatter.cpp b/sycl/test/esimd/gather_scatter.cpp index d98b1a268ebf5..71ec0caa1bdf4 100644 --- a/sycl/test/esimd/gather_scatter.cpp +++ b/sycl/test/esimd/gather_scatter.cpp @@ -9,8 +9,8 @@ using namespace sycl::ext::intel::experimental::esimd; using namespace cl::sycl; -void kernel(accessor &buf) +void kernel( + accessor &buf) __attribute__((sycl_device)) { simd offsets(0, 1); simd v1(0, 1); @@ -22,8 +22,8 @@ void kernel(accessor(buf.get_pointer(), v0, offsets); } -void kernel(accessor &buf) +void kernel( + accessor &buf) __attribute__((sycl_device)) { simd offsets(0, 1); simd v1(0, 1); @@ -38,7 +38,7 @@ void kernel(accessor &buf) + access::target::device> &buf) __attribute__((sycl_device)) { simd offsets(0, 1); simd v1(0, 1); diff --git a/sycl/test/esimd/gather_scatter_rgba.cpp b/sycl/test/esimd/gather_scatter_rgba.cpp index 2fc8ba3a965e1..2b927c958d228 100644 --- a/sycl/test/esimd/gather_scatter_rgba.cpp +++ b/sycl/test/esimd/gather_scatter_rgba.cpp @@ -11,8 +11,8 @@ using namespace sycl::ext::intel::experimental::esimd; using namespace cl::sycl; -void kernel(accessor &buf) SYCL_ESIMD_FUNCTION { +void kernel(accessor + &buf) SYCL_ESIMD_FUNCTION { simd offsets(0, 1); simd v1(0, 1); diff --git a/sycl/test/esimd/intrins_trans.cpp b/sycl/test/esimd/intrins_trans.cpp index 989ae37b6818b..563bc65b8beac 100644 --- a/sycl/test/esimd/intrins_trans.cpp +++ b/sycl/test/esimd/intrins_trans.cpp @@ -117,7 +117,7 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd foo() { { sycl::accessor + sycl::access::target::device> acc; simd offsets = 1; simd pred{1, 0, 1, 0, 1, 0, 1, 0}; diff --git a/sycl/test/esimd/simd_copy_to_copy_from.cpp b/sycl/test/esimd/simd_copy_to_copy_from.cpp index 61e4f8ee32c83..2bd09b588631c 100644 --- a/sycl/test/esimd/simd_copy_to_copy_from.cpp +++ b/sycl/test/esimd/simd_copy_to_copy_from.cpp @@ -16,9 +16,9 @@ using namespace cl::sycl; // --- Postive tests. -SYCL_EXTERNAL void kernel1( - accessor - &buf) SYCL_ESIMD_FUNCTION { +SYCL_EXTERNAL void +kernel1(accessor &buf) + SYCL_ESIMD_FUNCTION { simd v1(0, 1); simd v0; v0.copy_from(buf, 0); @@ -50,8 +50,8 @@ kernel3(accessor &buf) } // Incompatible mode (write). -SYCL_EXTERNAL void kernel4( - accessor &buf) +SYCL_EXTERNAL void +kernel4(accessor &buf) SYCL_ESIMD_FUNCTION { simd v; // CHECK: simd_copy_to_copy_from.cpp:58{{.*}}error: no matching member function for call to 'copy_from' @@ -59,8 +59,8 @@ SYCL_EXTERNAL void kernel4( } // Incompatible mode (read). -SYCL_EXTERNAL void kernel5( - accessor &buf) +SYCL_EXTERNAL void +kernel5(accessor &buf) SYCL_ESIMD_FUNCTION { simd v(0, 1); // CHECK: simd_copy_to_copy_from.cpp:67{{.*}}error: no matching member function for call to 'copy_to' diff --git a/sycl/test/extensions/sub_group_as.cpp b/sycl/test/extensions/sub_group_as.cpp index 52e4630685613..f61831aef4a45 100644 --- a/sycl/test/extensions/sub_group_as.cpp +++ b/sycl/test/extensions/sub_group_as.cpp @@ -24,9 +24,8 @@ int main(int argc, char *argv[]) { { cl::sycl::buffer buf(host_mem, N); queue.submit([&](cl::sycl::handler &cgh) { - auto global = - buf.get_access(cgh); + auto global = buf.get_access(cgh); sycl::accessor local(N, cgh); diff --git a/sycl/test/multi_ptr/ctad.cpp b/sycl/test/multi_ptr/ctad.cpp index d66b3fbfc5adc..a37c94fda737a 100644 --- a/sycl/test/multi_ptr/ctad.cpp +++ b/sycl/test/multi_ptr/ctad.cpp @@ -15,15 +15,20 @@ int main() { using sycl::access::address_space; using sycl::access::mode; using sycl::access::target; + using deviceAcc = sycl::accessor; using globlAcc = sycl::accessor; using constAcc = sycl::accessor; using localAcc = sycl::accessor; + using deviceCTAD = decltype(sycl::multi_ptr(std::declval())); using globlCTAD = decltype(sycl::multi_ptr(std::declval())); using constCTAD = decltype(sycl::multi_ptr(std::declval())); using localCTAD = decltype(sycl::multi_ptr(std::declval())); + using deviceMPtr = sycl::multi_ptr; using globlMPtr = sycl::multi_ptr; using constMPtr = sycl::multi_ptr; using localMPtr = sycl::multi_ptr; + static_assert(std::is_same::value); + static_assert(std::is_same::value); static_assert(std::is_same::value); static_assert(std::is_same::value); static_assert(std::is_same::value);