Add cast_to function

Author: stijnh

docs/build_api.py

@@ -199,6 +199,7 @@ def build_index_page(groups):
             ("where", "where(const C&)"),
         ],
         "Memory read/write": [
+            "cast_to",
             ("load", "load(const T*, const I&)"),
             ("load", "load(const T*, const I&, const M&)"),
             ("loadn", "loadn(const T*, ptrdiff_t)"),

include/kernel_float/memory.h

@@ -223,6 +223,47 @@ KERNEL_FLOAT_INLINE void storen(const V& values, T* ptr, ptrdiff_t offset, ptrdi
 #define KERNEL_FLOAT_ASSUME_ALIGNED(ptr, alignment) (ptr)
 #endif
 
+template<typename T, size_t N>
+struct AssignConversionProxy {
+    KERNEL_FLOAT_INLINE
+    explicit AssignConversionProxy(T* ptr) : ptr_(ptr) {}
+
+    template<typename U>
+    KERNEL_FLOAT_INLINE AssignConversionProxy& operator=(U&& values) {
+        auto indices = range<ptrdiff_t, N>();
+        detail::store_impl<T, N>::call(
+            ptr_,
+            convert_storage<T, N>(std::forward<U>(values)).data(),
+            indices.data());
+
+        return *this;
+    }
+
+  private:
+    T* ptr_;
+};
+
+/**
+ * Takes a reference to a vector and returns a special proxy object that automatically performs the correct conversion
+ * when a vector of a different element type is assigned. This is useful to perform implicit type conversions.
+ *
+ * For example, let assume that a line like `x = expression;` would not compile since `x` and `expressions` are
+ * vectors of different element types. Then it is possible to use `cast_to(x) = expression;` to fix this error,
+ * which possibly introduces a type conversion.
+ *
+ * Example
+ * =======
+ * ```
+ * vec<float, 2> x;
+ * vec<double, 2> y = {1.0, 2.0};
+ * cast_to(x) = y;  // normally, the line `x = y;` would not compile, but `cast_to` make this possible
+ * ```
+ */
+template<typename T, typename E>
+KERNEL_FLOAT_INLINE AssignConversionProxy<T, E::value> cast_to(vector<T, E>& input) {
+    return AssignConversionProxy<T, E::value>(input.data());
+}
+
 /**
  * Represents a pointer of type ``T*`` that is guaranteed to be aligned to ``alignment`` bytes.
  */

include/kernel_float/tiling.h

@@ -127,4 +127,19 @@ struct store_test {
 };
 
 REGISTER_TEST_CASE("store", store_test, int, float, double)
-REGISTER_TEST_CASE_GPU("store", store_test, __half, __nv_bfloat16)
+REGISTER_TEST_CASE_GPU("store", store_test, __half, __nv_bfloat16)
+
+struct assign_conversion_test {
+    template<typename T, size_t... I, size_t N = sizeof...(I)>
+    __host__ __device__ void operator()(generator<T> gen, std::index_sequence<I...>) {
+        kf::vec<T, N> x = {gen.next(I)...};
+        kf::vec<float, N> y;
+
+        kf::cast_to(y) = x;
+
+        ASSERT_EQ_ALL(float(x[I]), y[I]);
+    }
+};
+
+REGISTER_TEST_CASE("assign conversion", assign_conversion_test, int, float, double)
+REGISTER_TEST_CASE_GPU("assign conversion", assign_conversion_test, __half, __nv_bfloat16)