From d5634578baae88bdeecec57927cce9e946050190 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Tue, 13 Jul 2021 16:38:27 -0700 Subject: [PATCH 1/2] [ESIMD] rename gather4/scatter4 to gather_rgba/scatter_rgba --- .../ExplicitSIMD/dpcpp-explicit-simd.md | 4 +- .../ext/intel/experimental/esimd/memory.hpp | 78 ++++++++++++++----- sycl/test/esimd/gather4_scatter4.cpp | 49 ------------ sycl/test/esimd/gather_scatter_rgba.cpp | 51 ++++++++++++ 4 files changed, 112 insertions(+), 70 deletions(-) delete mode 100644 sycl/test/esimd/gather4_scatter4.cpp create mode 100644 sycl/test/esimd/gather_scatter_rgba.cpp diff --git a/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md b/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md index 8c4134a096f33..27ff62c175aee 100644 --- a/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md +++ b/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md @@ -357,13 +357,13 @@ template typename std::enable_if<(n == 16 || n == 32), simd>::type - flat_load4(T *p, simd offsets, simd pred = 1); + gather_rgba(T *p, simd offsets, simd pred = 1); template typename std::enable_if<(n == 16 || n == 32), void>::type - flat_store4(T *p, simd vals, + scatter_rgba(T *p, simd vals, simd offsets, simd pred = 1); ``` diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp index dadcc437d0e5d..02cfd986a0410 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp @@ -373,42 +373,82 @@ ESIMD_INLINE ESIMD_NODEBUG void scalar_store(AccessorTy acc, uint32_t offset, scatter(acc, simd{val}, simd{offset}); } -// TODO @jasonsewall-intel -// Don't use '4' in the name - instead either make it a parameter or -// (if it must be constant) - try to deduce from other arguments. -// +/// Gathering read for the given starting pointer \p p and \p offsets. +/// Up to 4 data elements may be accessed at each address depending on the +/// enabled channel \p Mask. +/// \tparam T element type of the returned vector. Must be 4-byte. +/// \tparam N size of the \p offsets vector. Must be 16 or 32. +/// \tparam Mask represents a pixel's channel mask. +/// @param p the USM pointer. +/// @param offsets byte-offsets within the \p buffer to be gathered. +/// @param pred predication control used for masking lanes. +/// \ingroup sycl_esimd +template +ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< + (N == 16 || N == 32) && (sizeof(T) == 4), + simd> +gather_rgba(T *p, simd offsets, simd pred = 1) { + + simd offsets_i = convert(offsets); + simd addrs(reinterpret_cast(p)); + addrs = addrs + offsets_i; + return __esimd_flat_read4(addrs.data(), pred.data()); +} + /// Flat-address gather4. /// Only allow simd-16 and simd-32. /// \ingroup sycl_esimd template +__SYCL_DEPRECATED("use gather_rgba.") ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< (n == 16 || n == 32) && (sizeof(T) == 4), - simd> -gather4(T *p, simd offsets, simd pred = 1) { - - simd offsets_i = convert(offsets); - simd addrs(reinterpret_cast(p)); - addrs = addrs + offsets_i; - return __esimd_flat_read4(addrs.data(), pred.data()); + simd> gather4(T *p, + simd + offsets, + simd + pred = 1) { + return gather_rgba(p, offsets, pred); } -/// Flat-address scatter4. +/// Scatter write for the given starting pointer \p p and \p offsets. +/// Up to 4 data elements may be written at each address depending on the +/// enabled channel \p Mask. +/// \tparam T element type of the input vector. Must be 4-byte. +/// \tparam N size of the \p offsets vector. Must be 16 or 32. +/// \tparam Mask represents a pixel's channel mask. +/// @param p the USM pointer. +/// @param vals values to be written. +/// @param offsets byte-offsets within the \p buffer to be written. +/// @param pred predication control used for masking lanes. /// \ingroup sycl_esimd -template ESIMD_INLINE ESIMD_NODEBUG - typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) == 4), + typename sycl::detail::enable_if_t<(N == 16 || N == 32) && (sizeof(T) == 4), void> - scatter4(T *p, simd vals, - simd offsets, simd pred = 1) { - simd offsets_i = convert(offsets); - simd addrs(reinterpret_cast(p)); + scatter_rgba(T *p, simd vals, + simd offsets, simd pred = 1) { + simd offsets_i = convert(offsets); + simd addrs(reinterpret_cast(p)); addrs = addrs + offsets_i; - __esimd_flat_write4(addrs.data(), vals.data(), + __esimd_flat_write4(addrs.data(), vals.data(), pred.data()); } +/// Flat-address scatter4. +/// \ingroup sycl_esimd +template +__SYCL_DEPRECATED("use scatter_rgba.") +ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< + (n == 16 || n == 32) && (sizeof(T) == 4), + void> scatter4(T *p, simd vals, + simd offsets, simd pred = 1) { + scatter_rgba(p, vals, offsets, pred); +} + namespace detail { /// Check the legality of an atomic call in terms of size and type. /// \ingroup sycl_esimd diff --git a/sycl/test/esimd/gather4_scatter4.cpp b/sycl/test/esimd/gather4_scatter4.cpp deleted file mode 100644 index 6dd35fca26b35..0000000000000 --- a/sycl/test/esimd/gather4_scatter4.cpp +++ /dev/null @@ -1,49 +0,0 @@ -// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:" - -// This test checks compilation of ESIMD slm gather4/scatter4 APIs. Those which -// are deprecated must produce deprecation messages. - -#include -#include -#include -#include - -using namespace sycl::ext::intel::experimental::esimd; -using namespace cl::sycl; - -void kernel(accessor &buf) SYCL_ESIMD_FUNCTION { - simd offsets(0, 1); - simd v1(0, 1); - - // CHECK: gather4_scatter4.cpp:21{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated - // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: - auto v0 = gather4(buf.get_pointer(), offsets); - // CHECK: gather4_scatter4.cpp:24{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated - // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: - v0 = gather4(buf.get_pointer(), - offsets); - v0 = gather4(buf.get_pointer(), offsets); - - v0 = v0 + v1; - - // CHECK: gather4_scatter4.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated - // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: - scatter4(buf.get_pointer(), v0, offsets); - // CHECK: gather4_scatter4.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated - // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: - scatter4(buf.get_pointer(), v0, - offsets); - scatter4(buf.get_pointer(), v0, offsets); -} - -// A "border" between host and device compilations -// CHECK-LABEL: 4 warnings generated -// CHECK: gather4_scatter4.cpp:21{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated -// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: -// CHECK: gather4_scatter4.cpp:24{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated -// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: -// CHECK: gather4_scatter4.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated -// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: -// CHECK: gather4_scatter4.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated -// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: diff --git a/sycl/test/esimd/gather_scatter_rgba.cpp b/sycl/test/esimd/gather_scatter_rgba.cpp new file mode 100644 index 0000000000000..b2eb56614d65e --- /dev/null +++ b/sycl/test/esimd/gather_scatter_rgba.cpp @@ -0,0 +1,51 @@ +// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:" + +// This test checks compilation of ESIMD slm gather_rgba/scatter_rgba APIs. +// Those which are deprecated must produce deprecation messages. + +#include +#include +#include +#include + +using namespace sycl::ext::intel::experimental::esimd; +using namespace cl::sycl; + +void kernel(accessor &buf) SYCL_ESIMD_FUNCTION { + simd offsets(0, 1); + simd v1(0, 1); + + // CHECK: gather_scatter_rgba.cpp:21{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated + // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: + auto v0 = gather_rgba(buf.get_pointer(), offsets); + // CHECK: gather_scatter_rgba.cpp:24{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated + // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: + v0 = gather_rgba( + buf.get_pointer(), offsets); + v0 = + gather_rgba(buf.get_pointer(), offsets); + + v0 = v0 + v1; + + // CHECK: gather_scatter_rgba.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated + // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: + scatter_rgba(buf.get_pointer(), v0, offsets); + // CHECK: gather_scatter_rgba.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated + // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: + scatter_rgba(buf.get_pointer(), + v0, offsets); + scatter_rgba(buf.get_pointer(), v0, + offsets); +} + +// A "border" between host and device compilations +// CHECK-LABEL: 4 warnings generated +// CHECK: gather_scatter_rgba.cpp:21{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated +// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: +// CHECK: gather_scatter_rgba.cpp:24{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated +// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: +// CHECK: gather_scatter_rgba.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated +// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: +// CHECK: gather_scatter_rgba.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated +// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: From 867d92961baa09d2c09e2481573ba241c12627c8 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Fri, 16 Jul 2021 11:05:10 -0700 Subject: [PATCH 2/2] Fixed test after clang-format --- sycl/test/esimd/gather_scatter_rgba.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/test/esimd/gather_scatter_rgba.cpp b/sycl/test/esimd/gather_scatter_rgba.cpp index b2eb56614d65e..f32851b34f5b6 100644 --- a/sycl/test/esimd/gather_scatter_rgba.cpp +++ b/sycl/test/esimd/gather_scatter_rgba.cpp @@ -28,10 +28,10 @@ void kernel(accessor(buf.get_pointer(), v0, offsets); - // CHECK: gather_scatter_rgba.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated + // CHECK: gather_scatter_rgba.cpp:36{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated // CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note: scatter_rgba(buf.get_pointer(), v0, offsets); @@ -45,7 +45,7 @@ void kernel(accessor