[STF] Support dynamically added dependencies in host_launch (#8174)

* [STF] Use opaque struct pointers for C API handles

Replace `typedef void*` with `typedef struct <name>_t* <name>` for all
opaque handle types in the STF C API. This provides type safety at the
C level, preventing accidental mixing of different handle types.

The naming convention (e.g. stf_ctx_handle_t / stf_ctx_handle) matches
the Cython declarations in the stf_c_api branch for consistency.

Made-with: Cursor

* Add add_deps(task_dep_untyped) to host_launch_scope

Allow C/Python bindings to build dependencies incrementally
without requiring compile-time type information.

Cherry-picked C++ parts from afde3ce.

Made-with: Cursor

* Add untyped host_launch with host_launch_deps handle

Introduce host_launch_deps, an opaque handle that provides indexed
access to dependency data and optional user data inside host callbacks.
Supports get<T>(index), size(), user_data() with proper lifetime
management via a destructor callback.

New host_launch_deps class with if-constexpr untyped dispatch path
in operator->*, set_user_data() on host_launch_scope, graph-safe
resource management. Tests cover stream_ctx and graph_ctx (12 tests).

Cherry-picked C++ parts from 9db3b86.

Made-with: Cursor

* Fix build errors in host_launch untyped path

- Add set_user_data forwarding to unified_scope in context.cuh
- Use explicit host_launch_deps{} instead of {} for nvcc pair deduction

Cherry-picked C++ parts from 6f6bf42.

Made-with: Cursor

* Fix untyped dispatch for generic lambdas in host_launch

Use a private canary type to detect generic lambdas: if Fun also
accepts the canary, it is generic and should use the typed path.

Cherry-picked from af970ac.

Made-with: Cursor

* Fix untyped dispatch: use std::conjunction for nvcc compatibility

nvcc eagerly instantiates generic-lambda bodies during is_invocable_v
checks, causing hard errors when the lambda body uses members that
don't exist on host_launch_deps (e.g. data_handle()).

Use std::conjunction to short-circuit: is_invocable<Fun, host_launch_deps&>
is only instantiated when sizeof...(Deps) == 0.

Cherry-picked from 5356c1f.

Made-with: Cursor

* clang-format

* Fix incomplete type error for logical_data_untyped in host_launch_scope

host_launch_deps holds a std::vector<logical_data_untyped>, but
logical_data_untyped was only forward-declared through transitive
includes. GCC 9's new_allocator::deallocate applies alignof to the
element type, which requires a complete type. Include logical_data.cuh
to provide the full definition.

Made-with: Cursor

* Add comments to user_data and host_launch_deps APIs

Made-with: Cursor

* Add host_launch C bindings for stream and graph contexts

Expose the untyped host_launch API through the C interface, allowing
C and Python callers to schedule host callbacks as task graph nodes
with full dependency tracking and optional user data.

Made-with: Cursor

* Use opaque struct pointers for host_launch handles

Apply the same opaque struct pointer pattern to stf_host_launch_handle
and stf_host_launch_deps_handle, and update the corresponding casts
in the implementation.

Made-with: Cursor

* Use a less ambiguous variable name for clarity

* Minor comment improvement

Author: caugonnet

c/experimental/stf/include/cccl/c/experimental/stf/stf.h

@@ -440,6 +440,30 @@ typedef struct stf_task_handle_t* stf_task_handle;
 
 typedef struct stf_cuda_kernel_handle_t* stf_cuda_kernel_handle;
 
+//!
+//! \brief Opaque handle for a host launch scope
+//!
+//! A host launch scope schedules a user-provided C callback on the host
+//! as a proper task graph node, with full dependency tracking.
+//! Created with stf_host_launch_create() and destroyed with stf_host_launch_destroy().
+
+typedef struct stf_host_launch_handle_t* stf_host_launch_handle;
+
+//!
+//! \brief Opaque handle for host launch dependency data
+//!
+//! Passed to the host callback at invocation time.  Provides indexed
+//! access to the data of each dependency and to optional user data.
+
+typedef struct stf_host_launch_deps_handle_t* stf_host_launch_deps_handle;
+
+//!
+//! \brief C callback type for host launch
+//!
+//! \param deps Opaque handle to dependency data and user data
+
+typedef void (*stf_host_callback_fn)(stf_host_launch_deps_handle deps);
+
 //! \}
 
 //! \defgroup Context Context Management
@@ -1277,6 +1301,105 @@ void stf_cuda_kernel_destroy(stf_cuda_kernel_handle k);
 
 //! \}
 
+//! \defgroup HostLaunch Host Launch
+//! \brief Schedule a host callback as a task graph node with dependency tracking
+//!
+//! \details
+//! Host launch provides a way to run arbitrary host-side functions as part of
+//! the task graph. Unlike generic tasks where the user manually launches work
+//! on a stream, host launch automatically schedules a C callback via
+//! `cudaLaunchHostFunc` (stream context) or `cudaGraphAddHostNode` (graph context).
+//!
+//! This is the untyped counterpart of the C++ `ctx.host_launch(deps...)->*lambda`
+//! construct, designed for use from C and Python bindings.
+//! \{
+
+//! \brief Create a host launch scope on a regular context
+//!
+//! \param ctx Context handle
+//! \param[out] h Pointer to receive host launch handle
+//!
+//! \see stf_host_launch_destroy()
+void stf_host_launch_create(stf_ctx_handle ctx, stf_host_launch_handle* h);
+
+//! \brief Add a dependency to a host launch scope
+//!
+//! \param h Host launch handle
+//! \param ld Logical data handle
+//! \param m Access mode (STF_READ, STF_WRITE, STF_RW)
+//!
+//! \see stf_task_add_dep()
+void stf_host_launch_add_dep(stf_host_launch_handle h, stf_logical_data_handle ld, stf_access_mode m);
+
+//! \brief Set the debug symbol for a host launch scope
+//!
+//! \param h Host launch handle
+//! \param symbol Null-terminated string
+void stf_host_launch_set_symbol(stf_host_launch_handle h, const char* symbol);
+
+//! \brief Copy user data into the host launch scope
+//!
+//! The data is copied and later accessible via
+//! stf_host_launch_deps_get_user_data() inside the callback.
+//! An optional destructor is called on the copied buffer when the
+//! dependency handle is destroyed.
+//!
+//! \param h Host launch handle
+//! \param data Pointer to user data
+//! \param size Size of user data in bytes
+//! \param dtor Optional destructor for the copied data (may be NULL)
+void stf_host_launch_set_user_data(stf_host_launch_handle h, const void* data, size_t size, void (*dtor)(void*));
+
+//! \brief Submit the host callback and finalize the scope
+//!
+//! After this call, the callback will be invoked on the host when all
+//! read/write dependencies are satisfied.  The callback receives an
+//! opaque deps handle for accessing dependency data and user data.
+//!
+//! \param h Host launch handle
+//! \param callback Function pointer invoked on the host
+//!
+//! \see stf_host_launch_create()
+void stf_host_launch_submit(stf_host_launch_handle h, stf_host_callback_fn callback);
+
+//! \brief Destroy a host launch handle
+//!
+//! \param h Host launch handle
+//!
+//! \see stf_host_launch_create()
+void stf_host_launch_destroy(stf_host_launch_handle h);
+
+//! \brief Get the raw data pointer for a dependency
+//!
+//! Returns the host-side pointer to the data of the dependency at \p index.
+//! The pointer is valid only during the callback execution.
+//!
+//! \param deps Dependency handle
+//! \param index Zero-based dependency index
+//! \return Pointer to the data (as `slice<char>` data handle)
+void* stf_host_launch_deps_get(stf_host_launch_deps_handle deps, size_t index);
+
+//! \brief Get the byte size of a dependency
+//!
+//! \param deps Dependency handle
+//! \param index Zero-based dependency index
+//! \return Size in bytes
+size_t stf_host_launch_deps_get_size(stf_host_launch_deps_handle deps, size_t index);
+
+//! \brief Get the number of dependencies
+//!
+//! \param deps Dependency handle
+//! \return Number of dependencies
+size_t stf_host_launch_deps_size(stf_host_launch_deps_handle deps);
+
+//! \brief Get the user data pointer
+//!
+//! \param deps Dependency handle
+//! \return Pointer to the copied user data, or NULL if none was set
+void* stf_host_launch_deps_get_user_data(stf_host_launch_deps_handle deps);
+
+//! \}
+
 #ifdef __cplusplus
 }
 #endif

c/experimental/stf/src/stf.cu

@@ -411,6 +411,99 @@ void stf_cuda_kernel_destroy(stf_cuda_kernel_handle t)
   delete kernel_ptr;
 }
 
+// -----------------------------------------------------------------------------
+// Host launch
+// -----------------------------------------------------------------------------
+
+using host_launch_type = decltype(::std::declval<context>().host_launch());
+
+void stf_host_launch_create(stf_ctx_handle ctx, stf_host_launch_handle* h)
+{
+  _CCCL_ASSERT(ctx != nullptr, "context handle must not be null");
+  _CCCL_ASSERT(h != nullptr, "host launch handle output pointer must not be null");
+
+  auto* context_ptr = reinterpret_cast<context*>(ctx);
+  *h                = reinterpret_cast<stf_host_launch_handle>(new host_launch_type{context_ptr->host_launch()});
+}
+
+void stf_host_launch_add_dep(stf_host_launch_handle h, stf_logical_data_handle ld, stf_access_mode m)
+{
+  _CCCL_ASSERT(h != nullptr, "host launch handle must not be null");
+  _CCCL_ASSERT(ld != nullptr, "logical data handle must not be null");
+
+  auto* scope_ptr = reinterpret_cast<host_launch_type*>(h);
+  auto* ld_ptr    = reinterpret_cast<logical_data_untyped*>(ld);
+  scope_ptr->add_deps(task_dep_untyped(*ld_ptr, access_mode(m)));
+}
+
+void stf_host_launch_set_symbol(stf_host_launch_handle h, const char* symbol)
+{
+  _CCCL_ASSERT(h != nullptr, "host launch handle must not be null");
+  _CCCL_ASSERT(symbol != nullptr, "symbol must not be null");
+
+  auto* scope_ptr = reinterpret_cast<host_launch_type*>(h);
+  scope_ptr->set_symbol(symbol);
+}
+
+void stf_host_launch_set_user_data(stf_host_launch_handle h, const void* data, size_t size, void (*dtor)(void*))
+{
+  _CCCL_ASSERT(h != nullptr, "host launch handle must not be null");
+
+  auto* scope_ptr = reinterpret_cast<host_launch_type*>(h);
+  scope_ptr->set_user_data(data, size, dtor);
+}
+
+void stf_host_launch_submit(stf_host_launch_handle h, stf_host_callback_fn callback)
+{
+  _CCCL_ASSERT(h != nullptr, "host launch handle must not be null");
+  _CCCL_ASSERT(callback != nullptr, "callback must not be null");
+
+  auto* scope_ptr = reinterpret_cast<host_launch_type*>(h);
+  (*scope_ptr)->*[callback](reserved::host_launch_deps& deps) {
+    callback(reinterpret_cast<stf_host_launch_deps_handle>(&deps));
+  };
+}
+
+void stf_host_launch_destroy(stf_host_launch_handle h)
+{
+  _CCCL_ASSERT(h != nullptr, "host launch handle must not be null");
+
+  auto* scope_ptr = reinterpret_cast<host_launch_type*>(h);
+  delete scope_ptr;
+}
+
+void* stf_host_launch_deps_get(stf_host_launch_deps_handle deps, size_t index)
+{
+  _CCCL_ASSERT(deps != nullptr, "deps handle must not be null");
+
+  auto* d = reinterpret_cast<reserved::host_launch_deps*>(deps);
+  return d->get<slice<char>>(index).data_handle();
+}
+
+size_t stf_host_launch_deps_get_size(stf_host_launch_deps_handle deps, size_t index)
+{
+  _CCCL_ASSERT(deps != nullptr, "deps handle must not be null");
+
+  auto* d = reinterpret_cast<reserved::host_launch_deps*>(deps);
+  return d->get<slice<char>>(index).extent(0);
+}
+
+size_t stf_host_launch_deps_size(stf_host_launch_deps_handle deps)
+{
+  _CCCL_ASSERT(deps != nullptr, "deps handle must not be null");
+
+  auto* d = reinterpret_cast<reserved::host_launch_deps*>(deps);
+  return d->size();
+}
+
+void* stf_host_launch_deps_get_user_data(stf_host_launch_deps_handle deps)
+{
+  _CCCL_ASSERT(deps != nullptr, "deps handle must not be null");
+
+  auto* d = reinterpret_cast<reserved::host_launch_deps*>(deps);
+  return d->user_data();
+}
+
 // -----------------------------------------------------------------------------
 // Composite data place and execution place grid (for Python/cuTile multi-stream)
 // -----------------------------------------------------------------------------

c/experimental/stf/test/test_host_launch.cu

@@ -0,0 +1,159 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda_runtime.h>
+
+#include <c2h/catch2_test_helper.h>
+#include <cccl/c/experimental/stf/stf.h>
+
+__global__ void fill_kernel(int cnt, double* data, double value)
+{
+  int tid      = blockIdx.x * blockDim.x + threadIdx.x;
+  int nthreads = gridDim.x * blockDim.x;
+
+  for (int i = tid; i < cnt; i += nthreads)
+  {
+    data[i] = value + i;
+  }
+}
+
+struct verify_args
+{
+  size_t N;
+  bool* passed;
+};
+
+static void verify_callback(stf_host_launch_deps_handle deps)
+{
+  auto* v = static_cast<verify_args*>(stf_host_launch_deps_get_user_data(deps));
+
+  if (stf_host_launch_deps_size(deps) != 1)
+  {
+    *v->passed = false;
+    return;
+  }
+
+  if (stf_host_launch_deps_get_size(deps, 0) != v->N * sizeof(double))
+  {
+    *v->passed = false;
+    return;
+  }
+
+  auto* data = static_cast<double*>(stf_host_launch_deps_get(deps, 0));
+  for (size_t i = 0; i < v->N; i++)
+  {
+    if (fabs(data[i] - (42.0 + i)) > 1e-10)
+    {
+      *v->passed = false;
+      return;
+    }
+  }
+  *v->passed = true;
+}
+
+C2H_TEST("host_launch with stream context", "[host_launch]")
+{
+  const size_t N = 1024;
+
+  stf_ctx_handle ctx;
+  stf_ctx_create(&ctx);
+
+  double* host_data;
+  cudaMallocHost(&host_data, N * sizeof(double));
+  for (size_t i = 0; i < N; i++)
+  {
+    host_data[i] = 0.0;
+  }
+
+  stf_logical_data_handle lData;
+  stf_logical_data(ctx, &lData, host_data, N * sizeof(double));
+  stf_logical_data_set_symbol(lData, "data");
+
+  // Fill data via a kernel task
+  stf_task_handle t;
+  stf_task_create(ctx, &t);
+  stf_task_set_symbol(t, "fill");
+  stf_task_add_dep(t, lData, STF_WRITE);
+  stf_task_start(t);
+  double* dData = (double*) stf_task_get(t, 0);
+  fill_kernel<<<2, 128, 0, (cudaStream_t) stf_task_get_custream(t)>>>((int) N, dData, 42.0);
+  stf_task_end(t);
+  stf_task_destroy(t);
+
+  // Use host_launch to verify data on the host
+  bool passed = false;
+  verify_args vargs{N, &passed};
+
+  stf_host_launch_handle h;
+  stf_host_launch_create(ctx, &h);
+  stf_host_launch_set_symbol(h, "verify");
+  stf_host_launch_add_dep(h, lData, STF_READ);
+  stf_host_launch_set_user_data(h, &vargs, sizeof(vargs), nullptr);
+  stf_host_launch_submit(h, verify_callback);
+  stf_host_launch_destroy(h);
+
+  stf_logical_data_destroy(lData);
+  stf_ctx_finalize(ctx);
+
+  REQUIRE(passed);
+
+  cudaFreeHost(host_data);
+}
+
+C2H_TEST("host_launch with graph context", "[host_launch]")
+{
+  const size_t N = 1024;
+
+  stf_ctx_handle ctx;
+  stf_ctx_create_graph(&ctx);
+
+  double* host_data;
+  cudaMallocHost(&host_data, N * sizeof(double));
+  for (size_t i = 0; i < N; i++)
+  {
+    host_data[i] = 0.0;
+  }
+
+  stf_logical_data_handle lData;
+  stf_logical_data(ctx, &lData, host_data, N * sizeof(double));
+  stf_logical_data_set_symbol(lData, "data");
+
+  // Fill data via a generic task with stream capture
+  stf_task_handle t;
+  stf_task_create(ctx, &t);
+  stf_task_set_symbol(t, "fill");
+  stf_task_add_dep(t, lData, STF_WRITE);
+  stf_task_enable_capture(t);
+  stf_task_start(t);
+  double* dData       = (double*) stf_task_get(t, 0);
+  cudaStream_t stream = (cudaStream_t) stf_task_get_custream(t);
+  fill_kernel<<<2, 128, 0, stream>>>((int) N, dData, 42.0);
+  stf_task_end(t);
+  stf_task_destroy(t);
+
+  // Use host_launch to verify data on the host
+  bool passed = false;
+  verify_args vargs{N, &passed};
+
+  stf_host_launch_handle h;
+  stf_host_launch_create(ctx, &h);
+  stf_host_launch_set_symbol(h, "verify");
+  stf_host_launch_add_dep(h, lData, STF_READ);
+  stf_host_launch_set_user_data(h, &vargs, sizeof(vargs), nullptr);
+  stf_host_launch_submit(h, verify_callback);
+  stf_host_launch_destroy(h);
+
+  stf_logical_data_destroy(lData);
+  stf_ctx_finalize(ctx);
+
+  REQUIRE(passed);
+
+  cudaFreeHost(host_data);
+}