Rename wrap_ptr and assert_aligned to make_vec_ptr

Author: stijnh

examples/vector_add/main.cu

@@ -54,9 +54,9 @@ void run_kernel(int n) {
     int grid_size = (n + items_per_block - 1) / items_per_block;
     my_kernel<items_per_thread><<<grid_size, block_size>>>(
         n,
-        kf::assert_aligned(input_dev),
+        kf::make_vec_ptr(input_dev),
         constant,
-        kf::assert_aligned(output_dev));
+        kf::make_vec_ptr(output_dev));
 
     // Copy results back
     cuda_check(cudaMemcpy(output_dev, output_result.data(), sizeof(float) * n, cudaMemcpyDefault));

include/kernel_float/memory.h

@@ -119,7 +119,7 @@ constexpr size_t gcd(size_t a, size_t b) {
  * Returns true if a pointer having alignment of `a` bytes also has an alignment of `b` bytes. Returns false otherwise.
  */
 KERNEL_FLOAT_INLINE
-constexpr size_t alignment_divisible(size_t a, size_t b) {
+constexpr bool is_alignment_divisible(size_t a, size_t b) {
     return gcd(a, KERNEL_FLOAT_MAX_ALIGNMENT) % gcd(b, KERNEL_FLOAT_MAX_ALIGNMENT) == 0;
 }
 
@@ -482,7 +482,7 @@ struct vector_ptr {
         typename T2,
         size_t N2,
         size_t A2,
-        enable_if_t<detail::alignment_divisible(A2, Alignment), int> = 0>
+        enable_if_t<detail::is_alignment_divisible(A2, Alignment), int> = 0>
     KERNEL_FLOAT_INLINE vector_ptr(vector_ptr<T2, N2, U, A2> p) : data_(p.get()) {}
 
     /**
@@ -567,14 +567,14 @@ struct vector_ptr<T, N, const U, Alignment> {
         typename T2,
         size_t N2,
         size_t A2,
-        enable_if_t<detail::alignment_divisible(A2, Alignment), int> = 0>
+        enable_if_t<detail::is_alignment_divisible(A2, Alignment), int> = 0>
     KERNEL_FLOAT_INLINE vector_ptr(vector_ptr<T2, N2, const U, A2> p) : data_(p.get()) {}
 
     template<
         typename T2,
         size_t N2,
         size_t A2,
-        enable_if_t<detail::alignment_divisible(A2, Alignment), int> = 0>
+        enable_if_t<detail::is_alignment_divisible(A2, Alignment), int> = 0>
     KERNEL_FLOAT_INLINE vector_ptr(vector_ptr<T2, N2, U, A2> p) : data_(p.get()) {}
 
     KERNEL_FLOAT_INLINE vector_ref<value_type, N, const U, Alignment> operator*() const {
@@ -621,7 +621,7 @@ template<
     size_t N,
     typename U,
     size_t A,
-    typename = enable_if_t<(N * sizeof(U)) % A == 0>>
+    typename = enable_if_t<detail::is_alignment_divisible(N * sizeof(U), A)>>
 KERNEL_FLOAT_INLINE vector_ptr<T, N, U, A>& operator+=(vector_ptr<T, N, U, A>& p, size_t i) {
     return p = p + i;
 }
@@ -634,30 +634,30 @@ KERNEL_FLOAT_INLINE vector_ptr<T, N, U, A>& operator+=(vector_ptr<T, N, U, A>& p
  * @tparam U The type of the elements pointed to by the raw pointer.
  */
 template<typename T, size_t N = 1, typename U>
-KERNEL_FLOAT_INLINE vector_ptr<T, N, U> wrap_ptr(U* ptr) {
+KERNEL_FLOAT_INLINE vector_ptr<T, N, U> make_vec_ptr(U* ptr) {
     return vector_ptr<T, N, U> {ptr};
 }
 
+// Doxygen cannot deal with the `make_vec_ptr` being defined multiple times, we ignore the second definition.
+/// @cond IGNORE
 /**
  * Creates a `vector_ptr<T, N>` from a raw pointer `T*` by asserting a specific alignment `N`.
  *
  * @tparam N The alignment constraint for the vector_ptr.
  * @tparam T The type of the elements pointed to by the raw pointer.
  */
 template<size_t N, typename T>
-KERNEL_FLOAT_INLINE vector_ptr<T, N> assert_aligned(T* ptr) {
+KERNEL_FLOAT_INLINE vector_ptr<T, N> make_vec_ptr(T* ptr) {
     return vector_ptr<T, N> {ptr};
 }
 
-// Doxygen cannot deal with the `assert_aligned` being defined twice, we ignore the second definition.
-/// @cond IGNORE
 /**
  * Creates a `vector_ptr<T, 1>` from a raw pointer `T*`. The alignment is assumed to be KERNEL_FLOAT_MAX_ALIGNMENT.
  *
  * @tparam T The type of the elements pointed to by the raw pointer.
  */
-template<typename T>
-KERNEL_FLOAT_INLINE vector_ptr<T, 1, T, KERNEL_FLOAT_MAX_ALIGNMENT> assert_aligned(T* ptr) {
+template<decltype(nullptr) = nullptr, typename T>
+KERNEL_FLOAT_INLINE vector_ptr<T, 1, T, KERNEL_FLOAT_MAX_ALIGNMENT> make_vec_ptr(T* ptr) {
     return vector_ptr<T, 1, T, KERNEL_FLOAT_MAX_ALIGNMENT> {ptr};
 }
 /// @endcond

tests/memory.cu

@@ -156,7 +156,7 @@ struct vector_ptr_test {
         };
 
         {
-            kf::vector_ptr<const U, N> storage_ptr = kf::assert_aligned(storage.data);
+            kf::vector_ptr<const U, N> storage_ptr = kf::make_vec_ptr(storage.data);
             kf::vector_ptr<T, N, const U> ptr = storage_ptr;
             ASSERT_EQ(ptr.get(), static_cast<const U*>(storage.data));
 
@@ -175,7 +175,7 @@ struct vector_ptr_test {
         }
 
         {
-            kf::vector_ptr<U, N> storage_ptr = kf::assert_aligned(storage.data);
+            kf::vector_ptr<U, N> storage_ptr = kf::make_vec_ptr(storage.data);
             kf::vector_ptr<T, N, U> ptr = storage_ptr;
             ASSERT_EQ(ptr.get(), static_cast<U*>(storage.data));
 
@@ -237,18 +237,18 @@ struct vector_ptr_test {
         {
             U* ptr = nullptr;
 
-            auto a1 = kf::wrap_ptr<T>(ptr);
+            auto a1 = kf::make_vec_ptr<T>(ptr);
             ASSERT(std::is_same<decltype(a1), kf::vector_ptr<T, 1, U>>::value);
 
-            auto a2 = kf::wrap_ptr<T, 2>(ptr);
+            auto a2 = kf::make_vec_ptr<T, 2>(ptr);
             ASSERT(std::is_same<decltype(a2), kf::vector_ptr<T, 2, U>>::value);
 
-            auto a3 = kf::assert_aligned(ptr);
+            auto a3 = kf::make_vec_ptr(ptr);
             ASSERT(
                 std::is_same<decltype(a3), kf::vector_ptr<U, 1, U, KERNEL_FLOAT_MAX_ALIGNMENT>>::
                     value);
 
-            auto a4 = kf::assert_aligned<2>(ptr);
+            auto a4 = kf::make_vec_ptr<2>(ptr);
             ASSERT(std::is_same<decltype(a4), kf::vector_ptr<U, 2>>::value);
         }
     }