[STF] Add python bindings

.github/actions/workflow-build/build-workflow.py

@@ -1093,6 +1093,13 @@ def set_derived_tags(matrix_job):
             matrix_job["jobs"].remove(original_job)
             matrix_job["jobs"] += expanded_jobs
 
+    # Apply project-specific job exclusions (e.g. skip_jobs_on_windows when cxx is MSVC)
+    skip_on_windows = project.get("skip_jobs_on_windows")
+    if skip_on_windows and "cxx" in matrix_job and is_windows(matrix_job):
+        for job in skip_on_windows:
+            if job in matrix_job["jobs"]:
+                matrix_job["jobs"].remove(job)
+
 
 def next_explode_tag(matrix_job):
     non_exploded_tags = ["jobs"]

c/experimental/stf/CMakeLists.txt

@@ -1,6 +1,6 @@
 cmake_minimum_required(VERSION 3.21)
 
-project(CCCL_C_EXPERIMENTAL_STF LANGUAGES CUDA CXX C)
+project(CCCL_C_EXPERIMENTAL_STF LANGUAGES CUDA CXX)
 
 option(
   CCCL_C_EXPERIMENTAL_STF_ENABLE_TESTING

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

@@ -456,36 +456,87 @@ cudaStream_t stf_fence(stf_ctx_handle ctx);
 //!
 //! \brief Create logical data from existing memory buffer
 //!
-//! Creates logical data handle from an existing host memory buffer.
-//! STF takes ownership of data management during task execution.
+//! Creates logical data handle from existing memory buffer, assuming host data place.
+//! This is a convenience wrapper around stf_logical_data_with_place() with host placement.
 //!
 //! \param ctx Context handle
 //! \param[out] ld Pointer to receive logical data handle
-//! \param addr Pointer to existing data buffer
+//! \param addr Pointer to existing data buffer (assumed to be host memory)
 //! \param sz Size of data in bytes
 //!
 //! \pre ctx must be valid context handle
 //! \pre ld must not be NULL
-//! \pre addr must not be NULL
+//! \pre addr must not be NULL and point to host-accessible memory
 //! \pre sz must be greater than 0
 //! \post *ld contains valid logical data handle
 //!
-//! \note Original data pointer should not be accessed during task execution
-//! \note Data will be written back when logical data is destroyed or context finalized
+//! \note This function assumes host memory. For device/managed memory, use stf_logical_data_with_place()
+//! \note Equivalent to: stf_logical_data_with_place(ctx, ld, addr, sz, make_host_data_place())
 //!
 //! \par Example:
 //! \code
 //! float data[1024];
 //! stf_logical_data_handle ld;
-//! stf_logical_data(ctx, &ld, data, sizeof(data));
+//! stf_logical_data(ctx, &ld, data, sizeof(data));  // Assumes host memory
 //! // ... use in tasks ...
 //! stf_logical_data_destroy(ld);
 //! \endcode
 //!
-//! \see stf_logical_data_empty(), stf_logical_data_destroy()
+//! \see stf_logical_data_with_place(), stf_logical_data_empty(), stf_logical_data_destroy()
 
 void stf_logical_data(stf_ctx_handle ctx, stf_logical_data_handle* ld, void* addr, size_t sz);
 
+//!
+//! \brief Create logical data handle from address with data place specification
+//!
+//! Creates logical data handle from existing memory buffer, explicitly specifying where
+//! the memory is located (host, device, managed, etc.). This is the primary and recommended
+//! logical data creation function as it provides STF with essential memory location information
+//! for optimal data movement and placement strategies.
+//!
+//! \param ctx Context handle
+//! \param[out] ld Pointer to receive logical data handle
+//! \param addr Pointer to existing memory buffer
+//! \param sz Size of buffer in bytes
+//! \param dplace Data place specifying memory location
+//!
+//! \pre ctx must be valid context handle
+//! \pre ld must be valid pointer to logical data handle pointer
+//! \pre addr must point to valid memory of at least sz bytes
+//! \pre sz must be greater than 0
+//! \pre dplace must be valid data place (not invalid)
+//!
+//! \post *ld contains valid logical data handle on success
+//! \post Caller owns returned handle (must call stf_logical_data_destroy())
+//!
+//! \par Examples:
+//! \code
+//! // GPU device memory (recommended for CUDA arrays)
+//! float* device_ptr;
+//! cudaMalloc(&device_ptr, 1000 * sizeof(float));
+//! stf_data_place dplace = make_device_data_place(0);
+//! stf_logical_data_handle ld;
+//! stf_logical_data_with_place(ctx, &ld, device_ptr, 1000 * sizeof(float), dplace);
+//!
+//! // Host memory
+//! float* host_data = new float[1000];
+//! stf_data_place host_place = make_host_data_place();
+//! stf_logical_data_handle ld_host;
+//! stf_logical_data_with_place(ctx, &ld_host, host_data, 1000 * sizeof(float), host_place);
+//!
+//! // Managed memory
+//! float* managed_ptr;
+//! cudaMallocManaged(&managed_ptr, 1000 * sizeof(float));
+//! stf_data_place managed_place = make_managed_data_place();
+//! stf_logical_data_handle ld_managed;
+//! stf_logical_data_with_place(ctx, &ld_managed, managed_ptr, 1000 * sizeof(float), managed_place);
+//! \endcode
+//!
+//! \see make_device_data_place(), make_host_data_place(), make_managed_data_place()
+
+void stf_logical_data_with_place(
+  stf_ctx_handle ctx, stf_logical_data_handle* ld, void* addr, size_t sz, stf_data_place dplace);
+
 //!
 //! \brief Set symbolic name for logical data
 //!

c/experimental/stf/src/stf.cu

@@ -44,12 +44,44 @@ cudaStream_t stf_fence(stf_ctx_handle ctx)
 }
 
 void stf_logical_data(stf_ctx_handle ctx, stf_logical_data_handle* ld, void* addr, size_t sz)
+{
+  // Convenience wrapper: assume host memory
+  stf_logical_data_with_place(ctx, ld, addr, sz, make_host_data_place());
+}
+
+void stf_logical_data_with_place(
+  stf_ctx_handle ctx, stf_logical_data_handle* ld, void* addr, size_t sz, stf_data_place dplace)
 {
   assert(ctx);
   assert(ld);
 
   auto* context_ptr = static_cast<context*>(ctx);
-  auto ld_typed     = context_ptr->logical_data(make_slice((char*) addr, sz));
+
+  // Convert C data_place to C++ data_place
+  cuda::experimental::stf::data_place cpp_dplace;
+  switch (dplace.kind)
+  {
+    case STF_DATA_PLACE_HOST:
+      cpp_dplace = cuda::experimental::stf::data_place::host();
+      break;
+    case STF_DATA_PLACE_DEVICE:
+      cpp_dplace = cuda::experimental::stf::data_place::device(dplace.u.device.dev_id);
+      break;
+    case STF_DATA_PLACE_MANAGED:
+      cpp_dplace = cuda::experimental::stf::data_place::managed();
+      break;
+    case STF_DATA_PLACE_AFFINE:
+      cpp_dplace = cuda::experimental::stf::data_place::affine();
+      break;
+    default:
+      // Invalid data place - this should not happen with valid input
+      assert(false && "Invalid data_place kind");
+      cpp_dplace = cuda::experimental::stf::data_place::host(); // fallback
+      break;
+  }
+
+  // Create logical data with the specified data place
+  auto ld_typed = context_ptr->logical_data(make_slice((char*) addr, sz), cpp_dplace);
 
   // Store the logical_data_untyped directly as opaque pointer
   *ld = new logical_data_untyped{ld_typed};

ci/build_common.sh

@@ -12,7 +12,8 @@ cd "$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )";
 
 # Script defaults
 VERBOSE=${VERBOSE:-}
-HOST_COMPILER=${CXX:-g++} # $CXX if set, otherwise `g++`
+# Prefer CUDAHOSTCXX when set (e.g. cross-compilation) so -cmake-options and env agree
+HOST_COMPILER=${CUDAHOSTCXX:-${CXX:-g++}}
 CXX_STANDARD=17
 CUDA_COMPILER=${CUDACXX:-nvcc} # $CUDACXX if set, otherwise `nvcc`
 CUDA_ARCHS= # Empty, use presets by default.

ci/matrix.yaml

@@ -468,6 +468,7 @@ jobs:
   test_py_coop:     { name: "Test cuda.coop",  gpu: true,  needs: 'build_py_wheel', force_producer_ctk: "pybuild", invoke: { prefix: 'test_cuda_coop'} }
   test_py_par:      { name: "Test cuda.compute",     gpu: true,  needs: 'build_py_wheel', force_producer_ctk: "pybuild", invoke: { prefix: 'test_cuda_compute'} }
   test_py_examples: { name: "Test cuda.cccl.examples",     gpu: true,  needs: 'build_py_wheel', force_producer_ctk: "pybuild", invoke: { prefix: 'test_cuda_cccl_examples'} }
+  test_py_stf:      { name: "Test cuda.stf",               gpu: true,  needs: 'build_py_wheel', force_producer_ctk: "pybuild", invoke: { prefix: 'test_cuda_stf'} }
 
   # Run jobs for 'target' project (ci/util/build_and_test_targets.sh):
   run_cpu: { gpu: false }
@@ -524,7 +525,9 @@ projects:
     name: "Python"
     job_map:
       build: ['build_py_wheel']
-      test:  ['test_py_headers', 'test_py_coop', 'test_py_par', 'test_py_examples']
+      test:  ['test_py_headers', 'test_py_coop', 'test_py_par', 'test_py_examples', 'test_py_stf']
+    # STF C API and Python bindings are not built for MSVC
+    skip_jobs_on_windows: ['test_py_stf']
   cccl_c_parallel:
     name: 'CCCL C Parallel'
     stds: [20]

ci/test_cuda_stf_python.sh

@@ -0,0 +1,30 @@
+#!/bin/bash
+
+set -euo pipefail
+
+ci_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+source "$ci_dir/pyenv_helper.sh"
+
+# Parse common arguments
+source "$ci_dir/util/python/common_arg_parser.sh"
+parse_python_args "$@"
+cuda_major_version=$(nvcc --version | grep release | awk '{print $6}' | tr -d ',' | cut -d '.' -f 1 | cut -d 'V' -f 2)
+
+# Setup Python environment
+setup_python_env "${py_version}"
+
+# Fetch or build the cuda_cccl wheel:
+if [[ -n "${GITHUB_ACTIONS:-}" ]]; then
+  wheel_artifact_name=$("$ci_dir/util/workflow/get_wheel_artifact_name.sh")
+  "$ci_dir/util/artifacts/download.sh" ${wheel_artifact_name} /home/coder/cccl/
+else
+  "$ci_dir/build_cuda_cccl_python.sh" -py-version "${py_version}"
+fi
+
+# Install cuda_cccl
+CUDA_CCCL_WHEEL_PATH="$(ls /home/coder/cccl/wheelhouse/cuda_cccl-*.whl)"
+python -m pip install "${CUDA_CCCL_WHEEL_PATH}[test-cu${cuda_major_version}]"
+
+# Run tests for STF module
+cd "/home/coder/cccl/python/cuda_cccl/tests/"
+python -m pytest -n auto -v stf/

cudax/examples/stf/void_data_interface.cu

@@ -49,5 +49,9 @@ int main()
     return cuda_kernel_desc{dummy_kernel, 16, 128, 0};
   };
 
+  EXPECT(token.is_void_interface());
+  EXPECT(token2.is_void_interface());
+  EXPECT(token3.is_void_interface());
+
   ctx.finalize();
 }

cudax/test/stf/cpp/task_get_stream.cu

@@ -24,7 +24,8 @@ void test_stream()
   context ctx;
 
   auto token = ctx.token();
-  auto t     = ctx.task(token.write());
+  EXPECT(token.is_void_interface());
+  auto t = ctx.task(token.write());
   t.start();
   cudaStream_t s = t.get_stream();
   EXPECT(s != nullptr);

python/cuda_cccl/CMakeLists.txt

@@ -17,9 +17,34 @@ message(
 
 # Build cccl.c.parallel and add CCCL's install rules
 set(_cccl_root ../..)
+
+# Build and install C++ library first
 set(CCCL_TOPLEVEL_PROJECT ON) # Enable the developer builds
 set(CCCL_ENABLE_C_PARALLEL ON) # Build the cccl.c.parallel library
+# STF C interface and Python bindings are not compatible with MSVC; disable on Windows.
+if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
+  set(CCCL_ENABLE_C_EXPERIMENTAL_STF OFF)
+  message(STATUS "STF disabled: not compatible with MSVC")
+else()
+  set(CCCL_ENABLE_C_EXPERIMENTAL_STF ON)
+endif()
+set(CCCL_ENABLE_UNSTABLE ON) # Enable unstable features
+
+# Disable all testing, examples, and benchmarks - we only want the libraries
+set(CCCL_ENABLE_TESTING OFF)
+set(CCCL_ENABLE_EXAMPLES OFF)
+set(CCCL_ENABLE_BENCHMARKS OFF)
+set(CCCL_C_PARALLEL_ENABLE_TESTING OFF)
+set(CCCL_C_EXPERIMENTAL_STF_ENABLE_TESTING OFF)
+# Note: CCCL_ENABLE_CUDAX must be ON because STF depends on it (via CCCL_ENABLE_UNSTABLE)
+# But disable cudax tests, examples, and header testing
+set(cudax_ENABLE_TESTING OFF)
+set(cudax_ENABLE_EXAMPLES OFF)
+set(cudax_ENABLE_HEADER_TESTING OFF)
+set(cudax_ENABLE_CUFILE OFF)
 set(CCCL_C_PARALLEL_LIBRARY_OUTPUT_DIRECTORY ${SKBUILD_PROJECT_NAME})
+set(CCCL_C_EXPERIMENTAL_STF_LIBRARY_OUTPUT_DIRECTORY ${SKBUILD_PROJECT_NAME})
+
 # Just install the rest:
 set(libcudacxx_ENABLE_INSTALL_RULES ON)
 set(CUB_ENABLE_INSTALL_RULES ON)
@@ -34,10 +59,29 @@ add_subdirectory(${_cccl_root} _parent_cccl)
 set(CMAKE_INSTALL_LIBDIR "${old_libdir}") # pop
 set(CMAKE_INSTALL_INCLUDEDIR "${old_includedir}") # pop
 
-# ensure the destination directory exists
+# Create CCCL::cudax alias for STF (normally created by cccl-config.cmake)
+if (
+  CCCL_ENABLE_C_EXPERIMENTAL_STF
+  AND TARGET cudax::cudax
+  AND NOT TARGET CCCL::cudax
+)
+  add_library(CCCL::cudax ALIAS cudax::cudax)
+endif()
+
+# Ensure the destination directory exists
 file(MAKE_DIRECTORY "cuda/compute/${CUDA_VERSION_DIR}/cccl")
+if (CCCL_ENABLE_C_EXPERIMENTAL_STF)
+  file(MAKE_DIRECTORY "cuda/stf/${CUDA_VERSION_DIR}/cccl")
+endif()
 
 # Install version-specific binaries
+if (CCCL_ENABLE_C_EXPERIMENTAL_STF)
+  install(
+    TARGETS cccl.c.experimental.stf
+    DESTINATION cuda/stf/${CUDA_VERSION_DIR}/cccl
+  )
+endif()
+
 install(
   TARGETS cccl.c.parallel
   DESTINATION cuda/compute/${CUDA_VERSION_DIR}/cccl
@@ -116,6 +160,7 @@ python3_add_library(
   WITH_SOABI
   "${_generated_extension_src}"
 )
+
 add_dependencies(_bindings_impl cythonize_bindings_impl)
 target_link_libraries(
   _bindings_impl
@@ -126,3 +171,53 @@ target_link_libraries(
 set_target_properties(_bindings_impl PROPERTIES INSTALL_RPATH "$ORIGIN/cccl")
 
 install(TARGETS _bindings_impl DESTINATION cuda/compute/${CUDA_VERSION_DIR})
+
+if (CCCL_ENABLE_C_EXPERIMENTAL_STF)
+  message(STATUS "STF Using Cython ${CYTHON_VERSION}")
+  set(
+    stf_pyx_source_file
+    "${cuda_cccl_SOURCE_DIR}/cuda/stf/_stf_bindings_impl.pyx"
+  )
+  set(
+    _stf_generated_extension_src
+    "${cuda_cccl_BINARY_DIR}/_stf_bindings_impl.c"
+  )
+  set(_stf_depfile "${cuda_cccl_BINARY_DIR}/_stf_bindings_impl.c.dep")
+  add_custom_command(
+    OUTPUT "${_stf_generated_extension_src}"
+    COMMAND "${Python3_EXECUTABLE}" -m cython
+    ARGS
+      ${CYTHON_FLAGS_LIST} "${stf_pyx_source_file}" --output-file
+      ${_stf_generated_extension_src}
+    DEPENDS "${stf_pyx_source_file}"
+    DEPFILE "${_stf_depfile}"
+    COMMENT "Cythonizing ${stf_pyx_source_file} for CUDA ${CUDA_VERSION_MAJOR}"
+  )
+  set_source_files_properties(
+    "${_stf_generated_extension_src}"
+    PROPERTIES GENERATED TRUE
+  )
+  add_custom_target(
+    cythonize_stf_bindings_impl
+    ALL
+    DEPENDS "${_stf_generated_extension_src}"
+  )
+
+  python3_add_library(
+    _stf_bindings_impl
+    MODULE
+    WITH_SOABI
+    "${_stf_generated_extension_src}"
+  )
+  add_dependencies(_stf_bindings_impl cythonize_stf_bindings_impl)
+  target_link_libraries(
+    _stf_bindings_impl
+    PRIVATE cccl.c.experimental.stf CUDA::cuda_driver
+  )
+  set_target_properties(
+    _stf_bindings_impl
+    PROPERTIES INSTALL_RPATH "$ORIGIN/cccl"
+  )
+
+  install(TARGETS _stf_bindings_impl DESTINATION cuda/stf/${CUDA_VERSION_DIR})
+endif()