improve host/device accessors documentation (NVIDIA#4220)

Author: fbusato

docs/libcudacxx/extended_api/mdspan/host_device_accessor.rst

@@ -71,33 +71,35 @@ Features
 
 **Memory spaces**
 
-+--------------------+------------------+-------------------+
-|                    | Host memory      | Device memory     |
-+====================+==================+===================+
-| ``host_mdspan``    | Allowed          | *Compile error*   |
-+--------------------+------------------+-------------------+
-| ``device_mdspan``  | *Compile error*  | Allowed           |
-+--------------------+------------------+-------------------+
-| ``managed_mdspan`` | Allowed *****    | Allowed *****     |
-+--------------------+------------------+-------------------+
-
-***** the validity of the *managed* memory space is checked at run-time in debug mode.
+*Host*, *device*, and *managed* ``mdspan`` can be created and "sliced" (``cuda::std::submdspan``) on any memory spaces. However, accessing to a specific memory space is restricted to the respective *accessor* type.
+
++----------------------------------+------------------+-------------------+
+| ``mdspan`` / memory space access | Host memory      | Device memory     |
++==================================+==================+===================+
+| ``host_mdspan``                  | Allowed          | *Compile error*   |
++----------------------------------+------------------+-------------------+
+| ``device_mdspan``                | *Compile error*  | Allowed           |
++----------------------------------+------------------+-------------------+
+| ``managed_mdspan``               | Allowed *****    | Allowed *****     |
++----------------------------------+------------------+-------------------+
+
+***** the validity of the *managed* memory space is checked at run-time in debug mode (host-side).
 
 **Conversions**
 
 +-----------------------------+------------------+-------------------+---------------------+
 |                             | ``host_mdspan``  | ``device_mdspan`` | ``managed_mdspan``  |
 +=============================+==================+===================+=====================+
-| ``host_mdspan``             | Allowed *****    | *Compile error*   | *Compile error*     |
+| ``host_mdspan``             | Allowed          | *Compile error*   | *Compile error*     |
 +-----------------------------+------------------+-------------------+---------------------+
-| ``device_mdspan``           | *Compile error*  | Allowed *****     | *Compile error*     |
+| ``device_mdspan``           | *Compile error*  | Allowed           | *Compile error*     |
 +-----------------------------+------------------+-------------------+---------------------+
-| ``managed_mdspan``          | Allowed          | Allowed           | Allowed *****       |
+| ``managed_mdspan``          | Allowed          | Allowed           | Allowed             |
 +-----------------------------+------------------+-------------------+---------------------+
 | Other mdspan                | Allowed          | Allowed           | Allowed             |
 +-----------------------------+------------------+-------------------+---------------------+
 
-***** the conversion is ``explicit`` if the base accessors are not directly convertible.
+*Note:* the conversion is ``explicit`` if the base accessor is not directly convertible.
 
 Example 1
 ---------
@@ -128,19 +130,20 @@ Example 1
     }
 
     int main() {
-        void* d_ptr;
+        int* d_ptr;
         cudaMalloc(&d_ptr, 4 * sizeof(int));
-        int   h_ptr[4];
-        cuda::host_mdspan<int, dim>   h_md{h_ptr, 4};
-        cuda::device_mdspan<int, dim> d_md{d_ptr, 4};
+        int                 h_ptr[4];
+        cuda::host_mdspan   h_md{h_ptr};
+        cuda::device_mdspan d_md{d_ptr, 4};
         kernel_d<<<1, 1>>>(d_md);    // ok
         // kernel_d<<<1, 1>>>(h_md); // compile error
         host_function_h(h_md);       // ok
         host_function_d(h_md);       // compile error
         // host_function_m(h_md);    // compile error
+        cudaFree(d_ptr);
     }
 
-`See example 1 on Godbolt 🔗 <https://godbolt.org/z/hW9faqsGW>`_
+`See example 1 on Godbolt 🔗 <https://godbolt.org/z/fezxsbjaq>`_
 
 Example 2
 ---------
@@ -162,17 +165,18 @@ Example 2
     }
 
     int main() {
-        void* m_ptr;
+        int* m_ptr;
         cudaMallocManaged(&m_ptr, 4 * sizeof(int));
-        cuda::managed_mdspan<int, dim> m_md{m_ptr, 4};
+        cuda::managed_mdspan m_md{m_ptr, 4};
         kernel_d<<<1, 1>>>(m_md); // ok
         host_function_h(m_md);    // ok
 
-        cuda::managed_mdspan<int, dim> m_md2{d_ptr, 4};
+        cuda::managed_mdspan m_md2{d_ptr, 4};
         m_md2[0]; // run-time error
+        cudaFree(d_ptr);
     }
 
-`See example 2 on Godbolt 🔗 <https://godbolt.org/z/WxWfaas5h>`_
+`See example 2 on Godbolt 🔗 <https://godbolt.org/z/Kj39Pe4vP>`_
 
 
 Example 3
@@ -189,13 +193,13 @@ Conversion from other accessors:
     int main() {
         using cuda::std::layout_right;
         using cuda::std::aligned_accessor;
-        int                         h_ptr[4];
-        cuda::std::mdspan<int, dim> md{h_ptr, 4};
-        cuda::host_mdspan<int, dim> h_md = md; // ok
+        int               h_ptr[4];
+        cuda::std::mdspan md{h_ptr};
+        cuda::host_mdspan h_md = md; // ok
 
         cuda::std::mdspan<int, dim, layout_right, aligned_accessor<int, 8>> md_a{h_ptr, 4};
-        // cuda::host_mdspan<int, dim> h_md = md_a; // compile-error
-        cuda::host_mdspan<int, dim>    h_md{md_a};  // ok
+        // cuda::host_mdspan h_md = md_a; // compile-error
+        cuda::host_mdspan    h_md{md_a};  // ok
     }
 
-`See example 3 on Godbolt 🔗 <https://godbolt.org/z/ja89roofx>`_
+`See example 3 on Godbolt 🔗 <https://godbolt.org/z/7dq7vcTWP>`_