[WIP] [DO NOT MERGE] add std::span support

sycl/include/CL/sycl/libcxx_span.hpp

@@ -148,14 +148,32 @@ _LIBCPP_PUSH_MACROS
 
 //CP adjust namespace declaration
 //_LIBCPP_BEGIN_NAMESPACE_STD
-namespace std {
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
 
 //CP living dangerously
 #define _LIBCPP_ASSERT(x, m) ((void)0)
 
 //CP - we make span available even outside C++20
 //#if _LIBCPP_STD_VER > 17
 
+//CP
+using std::numeric_limits;
+using std::true_type;
+using std::false_type;
+using std::remove_cv_t;
+using std::enable_if;
+using std::enable_if_t;
+using std::array;
+using std::void_t;
+using std::nullptr_t;
+using std::is_array_v;
+using std::declval;
+using std::is_convertible_v;
+using std::remove_pointer_t;
+using std::reverse_iterator;
+using std::is_const_v;
+
 inline constexpr size_t dynamic_extent = numeric_limits<size_t>::max();
 template <typename _Tp, size_t _Extent = dynamic_extent> class span;
 
@@ -366,11 +384,11 @@ class _LIBCPP_TEMPLATE_VIS span {
     _LIBCPP_INLINE_VISIBILITY constexpr rev_iterator        rbegin() const noexcept { return rev_iterator(end()); }
     _LIBCPP_INLINE_VISIBILITY constexpr rev_iterator          rend() const noexcept { return rev_iterator(begin()); }
 
-    _LIBCPP_INLINE_VISIBILITY span<const byte, _Extent * sizeof(element_type)> __as_bytes() const noexcept
-    { return span<const byte, _Extent * sizeof(element_type)>{reinterpret_cast<const byte *>(data()), size_bytes()}; }
+    _LIBCPP_INLINE_VISIBILITY span<const std::byte, _Extent * sizeof(element_type)> __as_bytes() const noexcept
+    { return span<const std::byte, _Extent * sizeof(element_type)>{reinterpret_cast<const std::byte *>(data()), size_bytes()}; }
 
-    _LIBCPP_INLINE_VISIBILITY span<byte, _Extent * sizeof(element_type)> __as_writable_bytes() const noexcept
-    { return span<byte, _Extent * sizeof(element_type)>{reinterpret_cast<byte *>(data()), size_bytes()}; }
+    _LIBCPP_INLINE_VISIBILITY span<std::byte, _Extent * sizeof(element_type)> __as_writable_bytes() const noexcept
+    { return span<std::byte, _Extent * sizeof(element_type)>{reinterpret_cast<std::byte *>(data()), size_bytes()}; }
 
 private:
     pointer    __data;
@@ -529,11 +547,11 @@ class _LIBCPP_TEMPLATE_VIS span<_Tp, dynamic_extent> {
     _LIBCPP_INLINE_VISIBILITY constexpr rev_iterator        rbegin() const noexcept { return rev_iterator(end()); }
     _LIBCPP_INLINE_VISIBILITY constexpr rev_iterator          rend() const noexcept { return rev_iterator(begin()); }
 
-    _LIBCPP_INLINE_VISIBILITY span<const byte, dynamic_extent> __as_bytes() const noexcept
-    { return {reinterpret_cast<const byte *>(data()), size_bytes()}; }
+    _LIBCPP_INLINE_VISIBILITY span<const std::byte, dynamic_extent> __as_bytes() const noexcept
+    { return {reinterpret_cast<const std::byte *>(data()), size_bytes()}; }
 
-    _LIBCPP_INLINE_VISIBILITY span<byte, dynamic_extent> __as_writable_bytes() const noexcept
-    { return {reinterpret_cast<byte *>(data()), size_bytes()}; }
+    _LIBCPP_INLINE_VISIBILITY span<std::byte, dynamic_extent> __as_writable_bytes() const noexcept
+    { return {reinterpret_cast<std::byte *>(data()), size_bytes()}; }
 
 private:
     pointer   __data;
@@ -571,7 +589,8 @@ template<class _Container>
 
 //CP
 //_LIBCPP_END_NAMESPACE_STD
-} // namespace std
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
 
 _LIBCPP_POP_MACROS
 

sycl/test/on-device/span/span.cpp

@@ -0,0 +1,104 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+#include <numeric>
+
+using namespace cl::sycl;
+
+
+void testSpanCapture() {
+    // This test creates spans that are backed by USM.
+    // ensures they can be captured by device lambda
+    // and that read and write operations function correctly
+    // across capture.
+    queue Q;
+
+    constexpr long numReadTests = 2;
+    const range<1> NumberOfReadTestsRange(numReadTests);
+    buffer<int, 1> SpanRead(NumberOfReadTestsRange);
+
+    // span from a vector
+    // We will create a vector, backed by a USM allocator. And a span from that.
+    typedef usm_allocator<int, usm::alloc::shared> vec_alloc;
+    // Create allocator for device associated with q
+    vec_alloc myAlloc(Q);
+    // Create std vector with the allocator
+    std::vector<int, vec_alloc> vecUSM(4, myAlloc);
+    std::iota(vecUSM.begin(), vecUSM.end(), 1);
+    sycl::span<int> vecUSM_span{vecUSM};
+    vecUSM_span[0] += 100; // 101  modify first value using span affordance.
+
+     // span from USM memory
+    int* usm_data = malloc_shared<int>(4, Q);
+    sycl::span<int> usm_span(usm_data, 4);
+    std::iota(usm_span.begin(), usm_span.end(), 1);
+    usm_span[0] += 100; // 101 modify first value using span affordance.
+
+    event E = Q.submit([&](handler &cgh){
+        auto span_read_acc = SpanRead.get_access<access::mode::write>(cgh);
+        cgh.single_task<class hi>([=] () {
+            // read from the spans.
+            span_read_acc[0] = vecUSM_span[0];
+            span_read_acc[1] = usm_span[0];
+
+            // write to the spans
+            vecUSM_span[1] += 1000;
+            usm_span[1] += 1000;
+        });
+    });
+    E.wait();
+
+    //check out the read operations, should have gotten 101 from each
+    auto span_read_acc = SpanRead.get_access<access::mode::read>();
+    for(int i=0; i < numReadTests; i++){
+        assert(span_read_acc[i] == 101 && "read check should have gotten 100");
+    }
+
+    //were the spans successfully modified via write?
+    assert(vecUSM_span[1] == 1002 && "vecUSM_span write check should have gotten 1001");
+    assert(usm_span[1] == 1002 && "usm_span write check should have gotten 1001");
+
+}
+
+void set_all_span_values(sycl::span<int> container, int v){
+    for(auto &e : container) 
+        e = v;
+}
+
+void testSpanOnDevice(){
+    // this test creates a simple span on device, 
+    // passes it to a function that operates on it
+    // and ensures it worked correctly
+    queue Q;
+    constexpr long numReadTests = 4;
+    const range<1> NumberOfReadTestsRange(numReadTests);
+    buffer<int, 1> SpanRead(NumberOfReadTestsRange);
+
+    event E = Q.submit([&](handler &cgh){
+        auto span_read_acc = SpanRead.get_access<access::mode::write>(cgh);
+        cgh.single_task<class ha>([=] () {
+            // create a span on device, pass it to function that modifies it
+            // read values back out.
+            int a[]{1, 2, 3, 4};
+            sycl::span<int> a_span{a};
+            set_all_span_values(a_span, 10);
+            for(int i=0; i < numReadTests; i++)
+                span_read_acc[i] = a_span[i];
+        });
+    });
+    E.wait();
+
+    //check out the read operations, should have gotten 10 from each
+    auto span_read_acc = SpanRead.get_access<access::mode::read>();
+    for(int i=0; i < numReadTests; i++){
+        assert(span_read_acc[i] == 10 && "read check should have gotten 10");
+    }
+}
+
+int main(){
+    testSpanCapture();
+    testSpanOnDevice();
+}