Add documentation for annotated_ptr and related APIs

docs/Dockerfile

@@ -0,0 +1,3 @@
+FROM jekyll/jekyll:4.0
+COPY Gemfile /Gemfile
+RUN /bin/bash -l -c "bundle install"

docs/Gemfile

@@ -6,4 +6,5 @@ group :jekyll_plugins do
   gem "jekyll-default-layout"        # GitHub Pages.
   gem "jekyll-titles-from-headings"  # GitHub Pages.
   gem "jekyll-relative-links"        # GitHub Pages.
+  gem "jekyll-remote-theme"          # GitHub Pages.
 end

docs/_config.yml

@@ -13,7 +13,8 @@ search.heading_level: 4
 # just-the-docs ignores these filenames by default.
 include: [ "contributing.md", "code_of_conduct.md" ]
 
-plugins:
+plugins_dir:
+  - jekyll-remote-theme
   - jekyll-optional-front-matter # GitHub Pages.
   - jekyll-default-layout        # GitHub Pages.
   - jekyll-titles-from-headings  # GitHub Pages.
@@ -81,6 +82,14 @@ defaults:
       has_children: true
       has_toc: false
       nav_order: 4
+  -
+    scope:
+      path: extended_api/memory_access_properties.md
+    values:
+      parent: Extended API
+      has_children: true
+      has_toc: false
+      nav_order: 5
   -
     scope:
       path: releases/changelog.md

docs/extended_api.md

@@ -17,6 +17,7 @@ nav_order: 3
 
 {% include_relative extended_api/asynchronous_operations.md %}
 
+{% include_relative extended_api/memory_access_properties.md %}
 
 [Thread Scopes]: ./extended_api/thread_groups.md
 [Thread Groups]: ./extended_api/thread_scopes.md

docs/extended_api/asynchronous_operations/memcpy_async.md

@@ -28,16 +28,42 @@ Defined in header `<cuda/pipeline>`:
 ```cuda
 // (3)
 template <typename Shape, cuda::thread_scope Scope>
+__host__ __device__
 void cuda::memcpy_async(void* destination, void const* source, Shape size,
                         cuda::pipeline<Scope>& pipeline);
 
 // (4)
 template <typename Group, typename Shape, cuda::thread_scope Scope>
+__host__ __device__
 void cuda::memcpy_async(Group const& group,
-                        void * destination, void const* source, Shape size,
+                        void* destination, void const* source, Shape size,
                         cuda::pipeline<Scope>& pipeline);
 ```
 
+Defined in header `<cuda/annotated_ptr>`:
+
+```cuda
+// (5)
+template <typename Dst, typename Src, typename SrcProperty, typename Shape, typename Sync>
+__host__ __device__
+void memcpy_async(Dst* dst, cuda::annotated_ptr<Src, SrcProperty> src, Shape size, Sync& sync);
+
+// (6)
+template<typename Dst, typename DstProperty, typename Src, typename SrcProperty, typename Shape, typename Sync>
+__host__ __device__
+void memcpy_async(cuda::annotated_ptr<Dst, DstProperty> dst, cuda::annotated_ptr<Src, SrcProperty> src, Shape size, Sync& sync);
+
+// (7)
+template<typename Group, typename Dst, typename Src, typename SrcProperty, typename Shape, typename Sync>
+__host__ __device__
+void memcpy_async(Group const& group, Dst* dst, cuda::annotated_ptr<Src, SrcProperty> src, Shape size, Sync& sync);
+
+// (8)
+template<typename Group, typename Dst, typename DstProperty, typename Src, typename SrcProperty, typename Shape, typename Sync>
+__host__ __device__
+void memcpy_async(Group const& group, cuda::annotated_ptr<Dst, DstProperty> dst, cuda::annotated_ptr<Src, SrcProperty> src, Shape size, Sync& sync);
+```
+
 `cuda::memcpy_async` asynchronously copies `size` bytes from the memory
   location pointed to by `source` to the memory location pointed to by
   `destination`.
@@ -51,6 +77,8 @@ Both objects are reinterpreted as arrays of `unsigned char`.
    in the current thread.
 4. Binds the asynchronous copy completion to `cuda::pipeline` and cooperatively
    issues the copy across all threads in `group`.
+5. 5-8: convenience wrappers using `cuda::annotated_ptr` where `Sync` is 
+   either `cuda::barrier` or `cuda::pipeline`.
 
 ## Notes
 

docs/extended_api/memory_access_properties.md

@@ -0,0 +1,13 @@
+## Memory access properties
+
+| [`cuda::annotated_ptr`]             | Binds an access property to a pointer. `(class template)` <br/><br/> 1.6.0 / CUDA 11.5 |
+| [`cuda::access_property`]           | Represents a memory access property. `(class)` <br/><br/> 1.6.0 / CUDA 11.5 |
+| [`cuda::apply_access_property`]     | Applies access property to memory location. `(function template)` <br/><br/> 1.6.0 / CUDA 11.5 |
+| [`cuda::associate_access_property`] | Associates access property with raw pointer. `(function template)` <br/><br/> 1.6.0 / CUDA 11.5 |
+| [`cuda::discard_memory`]            | Writes indeterminate values to memory. `(function)` <br/><br/> 1.6.0 / CUDA 11.5 |
+
+[`cuda::annotated_ptr`]: {{ "extended_api/memory_access_properties/annotated_ptr.html" | relative_url }}
+[`cuda::access_property`]: {{ "extended_api/memory_access_properties/access_property.html" | relative_url }}
+[`cuda::associate_access_property`]: {{ "extended_api/memory_access_properties/associate_access_property.html" | relative_url }}
+[`cuda::apply_access_property`]: {{ "extended_api/memory_access_properties/apply_access_property.html" | relative_url }}
+[`cuda::discard_memory`]: {{ "extended_api/memory_access_properties/discard_memory.html" | relative_url }}