diff --git a/nvbench/cuda_stream.cuh b/nvbench/cuda_stream.cuh
index 1c71372..f8f4a52 100644
--- a/nvbench/cuda_stream.cuh
+++ b/nvbench/cuda_stream.cuh
@@ -27,9 +27,19 @@
 namespace nvbench
 {
 
-// RAII wrapper for a cudaStream_t.
+/**
+ * Manages and provides access to a CUDA stream.
+ *
+ * May be owning or non-owning. If the stream is owned, it will be freed with
+ * `cudaStreamDestroy` when the `cuda_stream`'s lifetime ends. Non-owning
+ * `cuda_stream`s are sometimes referred to as views.
+ */
 struct cuda_stream
 {
+  /**
+   * Constructs a cuda_stream that owns a new stream, created with
+   * `cudaStreamCreate`.
+   */
   cuda_stream()
       : m_stream{[]() {
                    cudaStream_t s;
@@ -39,6 +49,12 @@ struct cuda_stream
                  stream_deleter{true}}
   {}
 
+  /**
+   * Constructs a `cuda_stream` from an explicit cudaStream_t.
+   *
+   * @param owning If true, `cudaStreamCreate(stream)` will be called from this
+   * `cuda_stream`'s destructor.
+   */
   cuda_stream(cudaStream_t stream, bool owning)
       : m_stream{stream, stream_deleter{owning}}
   {}
@@ -51,9 +67,14 @@ struct cuda_stream
   cuda_stream(cuda_stream &&)                 = default;
   cuda_stream &operator=(cuda_stream &&) = default;
 
+  /**
+   * @return The `cudaStream_t` managed by this `cuda_stream`.
+   * @{
+   */
   operator cudaStream_t() const { return m_stream.get(); }
 
   cudaStream_t get_stream() const { return m_stream.get(); }
+  /**@}*/
 
 private:
   struct stream_deleter
diff --git a/nvbench/launch.cuh b/nvbench/launch.cuh
index fae3f0e..4b973f3 100644
--- a/nvbench/launch.cuh
+++ b/nvbench/launch.cuh
@@ -23,6 +23,18 @@
 namespace nvbench
 {
 
+/**
+ * Configuration object used to communicate with a `KernelLauncher`.
+ *
+ * The `KernelLauncher` passed into `nvbench::state::exec` is required to
+ * accept an `nvbench::launch` argument:
+ *
+ * ```cpp
+ * state.exec([](nvbench::launch &launch) {
+ *   kernel<<<M, N, 0, launch.get_stream()>>>();
+ * }
+ * ```
+ */
 struct launch
 {
   explicit launch(const nvbench::cuda_stream &stream)
@@ -35,12 +47,18 @@ struct launch
   launch &operator=(const launch &) = delete;
   launch &operator=(launch &&) = default;
 
+  /**
+   * @return a CUDA stream that all kernels and other stream-ordered CUDA work
+   * must use. This stream can be changed by the `KernelGenerator` using the
+   * `nvbench::state::set_cuda_stream` method.
+   */
   __forceinline__ const nvbench::cuda_stream &get_stream() const
   {
     return m_stream;
   };
 
 private:
+  // The stream is owned by the `nvbench::state` associated with this launch.
   const nvbench::cuda_stream &m_stream;
 };