Restore MSVC compat

Also refactor ops

CMakeLists.txt

@@ -58,7 +58,7 @@ add_library(simsycl
     include/simsycl/sycl/multi_ptr.hh
     include/simsycl/sycl/nd_item.hh
     include/simsycl/sycl/nd_range.hh
-    include/simsycl/sycl/ops.hh
+    include/simsycl/sycl/binary_ops.hh
     include/simsycl/sycl/platform.hh
     include/simsycl/sycl/property.hh
     include/simsycl/sycl/queue.hh

include/simsycl/detail/allocation.hh

@@ -7,14 +7,33 @@
 
 namespace simsycl::detail {
 
+// aligned allocation wrapper for cross-platform support
+// NOTE: returned pointers must be freed with aligned_free
+inline void *aligned_alloc(size_t alignment, size_t size) {
+#if defined(_MSC_VER)
+    return _aligned_malloc(size, alignment);
+#else
+    return std::aligned_alloc(alignment, size);
+#endif
+}
+
+inline void aligned_free(void *ptr) {
+#if defined(_MSC_VER)
+    _aligned_free(ptr);
+#else
+    std::free(ptr);
+#endif
+}
+
+
 // floats and doubles filled with this pattern show up as "-nan"
 inline constexpr std::byte uninitialized_memory_pattern = std::byte(0xff);
 
 class allocation {
   public:
     allocation() = default;
     allocation(const size_t size_bytes, const size_t alignment_bytes)
-        : m_size(size_bytes), m_alignment(alignment_bytes), m_ptr(std::aligned_alloc(alignment_bytes, size_bytes)) {
+        : m_size(size_bytes), m_alignment(alignment_bytes), m_ptr(aligned_alloc(alignment_bytes, size_bytes)) {
         memset(m_ptr, static_cast<int>(uninitialized_memory_pattern), size_bytes);
     }
 
@@ -46,7 +65,7 @@ class allocation {
 
     void reset() {
         if(m_ptr != nullptr) {
-            free(m_ptr);
+            aligned_free(m_ptr);
             m_size = 0;
             m_alignment = 1;
             m_ptr = nullptr;

include/simsycl/detail/group_operation_impl.hh

@@ -9,10 +9,10 @@
 
 #include "simsycl/detail/allocation.hh"
 #include "simsycl/detail/check.hh"
+#include "simsycl/sycl/binary_ops.hh"
 #include "simsycl/sycl/concepts.hh"
 #include "simsycl/sycl/enums.hh"
 #include "simsycl/sycl/group.hh"
-#include "simsycl/sycl/ops.hh"
 #include "simsycl/sycl/sub_group.hh"
 
 namespace simsycl::detail {

include/simsycl/sycl/accessor.hh

@@ -97,7 +97,7 @@ template <typename Accessor, typename DataT>
 class accessor_iterator<Accessor, DataT, 0> {
   public:
     using value_type = DataT;
-    using difference_type = ssize_t;
+    using difference_type = long int;
     using reference = value_type &;
     using pointer = value_type *;
     using iterator_category = std::forward_iterator_tag;
@@ -589,7 +589,8 @@ class accessor<DataT, 0, AccessMode, AccessTarget, IsPlaceholder> : public simsy
 template <typename DataT, int Dimensions>
 class local_accessor final : public simsycl::detail::property_interface {
   private:
-    using property_compatibility = detail::property_compatibility_with<local_accessor<DataT, Dimensions>, property::no_init>;
+    using property_compatibility
+        = detail::property_compatibility_with<local_accessor<DataT, Dimensions>, property::no_init>;
 
   public:
     using value_type = DataT;

include/simsycl/sycl/binary_ops.hh

@@ -0,0 +1,138 @@
+#pragma once
+
+#include "concepts.hh"
+#include "type_traits.hh"
+
+#include <functional>
+
+namespace simsycl::sycl {
+
+// arithmetic
+
+using std::plus;
+template <typename OpT, typename T>
+struct has_known_identity<plus<OpT>, T> : std::bool_constant<detail::is_arithmetic_v<T>> {};
+template <typename OpT, Arithmetic T>
+struct known_identity<plus<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value = T{};
+};
+
+using std::multiplies;
+template <typename OpT, typename T>
+struct has_known_identity<multiplies<OpT>, T> : std::bool_constant<detail::is_arithmetic_v<T>> {};
+template <typename OpT, Arithmetic T>
+struct known_identity<multiplies<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value = T{1};
+};
+
+// bitwise boolean logic
+
+using std::bit_and;
+template <typename OpT, typename T>
+struct has_known_identity<bit_and<OpT>, T> : std::bool_constant<std::is_integral_v<T>> {};
+template <typename OpT, std::integral T>
+struct known_identity<bit_and<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value = ~T{};
+};
+
+using std::bit_or;
+template <typename OpT, typename T>
+struct has_known_identity<bit_or<OpT>, T> : std::bool_constant<std::is_integral_v<T>> {};
+template <typename OpT, std::integral T>
+struct known_identity<bit_or<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value = T{};
+};
+
+using std::bit_xor;
+template <typename OpT, typename T>
+struct has_known_identity<bit_xor<OpT>, T> : std::bool_constant<std::is_integral_v<T>> {};
+template <typename OpT, std::integral T>
+struct known_identity<bit_xor<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value = T{};
+};
+
+// logic
+
+using std::logical_and;
+template <typename OpT, typename T>
+struct has_known_identity<logical_and<OpT>, T> : std::bool_constant<std::is_same_v<std::remove_cv_t<T>, bool>> {};
+template <typename OpT, Boolean T>
+struct known_identity<logical_and<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value = true;
+};
+
+using std::logical_or;
+template <typename OpT, typename T>
+struct has_known_identity<logical_or<OpT>, T> : std::bool_constant<std::is_same_v<std::remove_cv_t<T>, bool>> {};
+template <typename OpT, Boolean T>
+struct known_identity<logical_or<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value = false;
+};
+
+// min/max
+
+template <typename T = void>
+struct minimum {
+    T operator()(const T &x, const T &y) const { return x <= y ? x : y; }
+};
+template <>
+struct minimum<void> {
+    template <typename T, typename U>
+    decltype(auto) operator()(T &&x, U &&y) const {
+        return x < y ? std::forward<T>(x) : std::forward<U>(y);
+    }
+};
+template <typename OpT, typename T>
+struct has_known_identity<minimum<OpT>, T> : std::bool_constant<detail::is_arithmetic_v<T>> {};
+template <typename OpT, Arithmetic T>
+struct known_identity<minimum<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value
+        = std::is_floating_point_v<T> ? std::numeric_limits<T>::infinity() : std::numeric_limits<T>::max();
+};
+
+template <typename T = void>
+struct maximum {
+    T operator()(const T &x, const T &y) const { return x >= y ? x : y; }
+};
+template <>
+struct maximum<void> {
+    template <typename T, typename U>
+    decltype(auto) operator()(T &&x, U &&y) const {
+        return x > y ? std::forward<T>(x) : std::forward<U>(y);
+    }
+};
+template <typename OpT, typename T>
+struct has_known_identity<maximum<OpT>, T> : std::bool_constant<detail::is_arithmetic_v<T>> {};
+template <typename OpT, Arithmetic T>
+struct known_identity<maximum<OpT>, T> {
+    static constexpr std::remove_cv_t<T> value
+        = std::is_floating_point_v<T> ? -std::numeric_limits<T>::infinity() : std::numeric_limits<T>::lowest();
+};
+
+} // namespace simsycl::sycl
+
+namespace simsycl::detail {
+
+template <typename T>
+struct is_function_object<simsycl::sycl::plus<T>> : std::true_type {};
+template <typename T>
+struct is_function_object<simsycl::sycl::multiplies<T>> : std::true_type {};
+
+template <typename T>
+struct is_function_object<simsycl::sycl::bit_and<T>> : std::true_type {};
+template <typename T>
+struct is_function_object<simsycl::sycl::bit_or<T>> : std::true_type {};
+template <typename T>
+struct is_function_object<simsycl::sycl::bit_xor<T>> : std::true_type {};
+
+template <typename T>
+struct is_function_object<simsycl::sycl::logical_and<T>> : std::true_type {};
+template <typename T>
+struct is_function_object<simsycl::sycl::logical_or<T>> : std::true_type {};
+
+template <typename T>
+struct is_function_object<simsycl::sycl::minimum<T>> : std::true_type {};
+template <typename T>
+struct is_function_object<simsycl::sycl::maximum<T>> : std::true_type {};
+
+} // namespace simsycl::detail

include/simsycl/sycl/group_algorithms.hh

@@ -1,9 +1,9 @@
 #pragma once
 
+#include "binary_ops.hh"
 #include "concepts.hh"
 #include "group.hh"
 #include "group_functions.hh"
-#include "ops.hh"
 #include "sub_group.hh"
 
 #include "simsycl/detail/check.hh"