Hacktoberfest 2026: những issue maintainer đã đánh dấu cho tháng Mười, đang mở và phù hợp người mới. Xem issue Hacktoberfest

sequence: missing transition (?)

Đang mở
  • Có gì không ổn?
#2,344 3 bình luận 0 reaction 0 người được giao Xem trên GitHub

Maintainer thường phản hồi trong vòng 1 ngày

Chưa có ai nhận issue này.

Đánh giá

Độ khó
5/5
Thời gian dự kiến
Hơn một tuần
Mức phù hợp với người mới
12/100
Loại issue
Lỗi
Độ rõ ràng
Cần làm rõ
Mức độ hoạt động
Sôi nổi
Công nghệ
cpp
Lĩnh vực
backend

Hướng nghiên cứu

Start with __sequence.hpp around line 255, where __state_->__start_next() runs inside the previous sender's completion, and with the static_assert on __completion_domain_of_t in the reproducer. The open question is semantic: whether exec::sequence should accept senders from incompatible domains, insert transitions such as continues_on, or reject them. Done means a maintainer decision on that behavior, followed by a fix and a test covering the CUDA-to-thread-pool case.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.

Mô tả

We think exec::sequence should either inject transitions, or reject completely when senders are incompatible.

Consider the following code.

#include "stdexec/execution.hpp"
#include "nvexec/stream_context.cuh"
#include "exec/static_thread_pool.hpp"

struct Functor {
    //! Seems like clang requires the @c __host__ annotation, otherwise it does not compile.
    __device__ __host__ void operator()() const noexcept {}
};

int main() {
    nvexec::stream_context stream_ctx{};
    exec::static_thread_pool pool{1};

    auto sndr = exec::sequence(
        stdexec::schedule(stream_ctx.get_scheduler()) | stdexec::then(Functor{}),
        stdexec::schedule(pool.get_scheduler()) | stdexec::then(Functor{}));

    static_assert(std::same_as<
        stdexec::__completion_domain_of_t<stdexec::set_value_t, decltype(sndr), stdexec::env<>>,
        experimental::execution::_pool_::_static_thread_pool::domain>);

    stdexec::sync_wait(std::move(sndr));
}

The stream_ctx scheduler is never told that it should "come back" to the CPU side, hence the runtime issue pasted hereinbelow (with a Docker reproducer).

What happens in this very case is that the sequence receiver (trivially copyable) passed to the nvexec machinery points to a pure CPU __state_ variable, so there is no chance that __state_->__start_next() can succeed when run on the GPU.

This is because the nvexec scheduler isn't a CPU-compatible one, it needs to be told when to come back on the CPU.

This could be mitigated by verifying if the nvexec domain is "incompatible" with the domain customizing exec::sequence (in this case the static thread pool domain), in which case the above code could be simply rejected, end of the story.

But we feel like there is a more insidious issue with the semantics of exec::sequence. Our current understanding is that the next sender is connected from within the completion context of the previous sender. This makes it quite clear that the nvexec case cannot work.

Consider the following hypothetical snippet: is it supposed to work for any schd_A/schd_B combination?

stdexec::sync_wait(exec::sequence(
    stdexec::schedule(schd_A) | stdexec::then(...),
    stdexec::schedule(schd_B) | stdexec::then(...));

If the answer is yes, how can we make it work?

  1. Should we introduce transitions represented by continues_on and schedule_from, or
  2. Should we wrap problematic senders in stdexec::affine or something?
  3. Anything else...

Joint work with @maartenarnst.

CUDA Exception: Warp MMU Fault
The exception was triggered at PC 0x76a00b2a7f20  stdexec::__seq::__rcvr1<stdexec::__sync_wait::__receiver<>, stdexec::__sched_env<nv::execution::_strm::stream_scheduler, stdexec::__env::__fwd<stdexec::__sync_wait::__env const&> > >::set_value()  (__sequence.hpp:255)

Thread 1 "test_sequence" received signal CUDA_EXCEPTION_30, Warp MMU Fault.
[Switching focus to CUDA kernel 0, grid 2, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 1, lane 0]
0x000076a00b2a7f30 in stdexec::__seq::__rcvr1<stdexec::__sync_wait::__receiver<>, stdexec::__sched_env<nv::execution::_strm::stream_scheduler, stdexec::__env::__fwd<stdexec::__sync_wait::__env const&> > >::set_value() (this=0x76a01ffffda0)
    at /tmp/stdexec-792e1342f8c983e205d899f840038d6d57fc9bd1/include/stdexec/__detail/__sequence.hpp:255
255	        __state_->__start_next();
(cuda-gdb) A debugging session is active.

	Inferior 1 [process 101] will be killed.
The Dockerfile to reproduce the failure (click to expand). 
ARG CUDA_VERSION=13.4.1
ARG UBUNTU=ubuntu26.04

FROM nvidia/cuda:${CUDA_VERSION}-devel-${UBUNTU}

ENV DEBIAN_FRONTEND=noninteractive

RUN apt update && apt install -y --no-install-recommends \
        ca-certificates curl wget \
        git \
        gnupg \
        lsb-release \
        software-properties-common \
        make cmake \
    && rm -rf /var/lib/apt/lists/*

ARG CLANG_VERSION=23

RUN wget -qO /tmp/llvm.sh https://apt.llvm.org/llvm.sh \
    && chmod +x /tmp/llvm.sh \
    && /tmp/llvm.sh ${CLANG_VERSION} \
    && rm /tmp/llvm.sh

ARG CMAKE_PACKAGE_MANAGER_VERSION=0.43.1

ADD https://github.com/cpm-cmake/CPM.cmake/releases/download/v${CMAKE_PACKAGE_MANAGER_VERSION}/CPM.cmake \
    /opt/cpm-${CMAKE_PACKAGE_MANAGER_VERSION}/

ARG STDEXEC_SHA=792e1342f8c983e205d899f840038d6d57fc9bd1

ADD https://github.com/NVIDIA/stdexec/archive/${STDEXEC_SHA}.tar.gz /tmp/

RUN cd /tmp && tar -xzf ${STDEXEC_SHA}.tar.gz

WORKDIR /tmp/

COPY <<EOF CMakeLists.txt
cmake_minimum_required(VERSION 4.0.0 FATAL_ERROR)

project(test CXX CUDA)

include(/opt/cpm-${CMAKE_PACKAGE_MANAGER_VERSION}/CPM.cmake)

cpmaddpackage(
  NAME
    stdexec
  SOURCE_DIR
    "/tmp/stdexec-${STDEXEC_SHA}"
  OPTIONS
    "STDEXEC_BUILD_TESTS OFF"
    "STDEXEC_BUILD_EXAMPLES OFF"
    "STDEXEC_INSTALL ON"
)

find_package(CUDAToolkit REQUIRED)

add_executable(test_sequence test_sequence.cpp)
target_link_libraries(test_sequence PRIVATE stdexec CUDA::cudart)
set_source_files_properties(test_sequence.cpp PROPERTIES LANGUAGE CUDA)
EOF

COPY <<EOF test_sequence.cpp
#include "stdexec/execution.hpp"
#include "nvexec/stream_context.cuh"
#include "exec/static_thread_pool.hpp"

struct Functor {
    //! Seems like clang requires the @c __host__ annotation, otherwise it does not compile.
    __device__ __host__ void operator()() const noexcept {}
};

int main() {
    nvexec::stream_context stream_ctx{};
    exec::static_thread_pool pool{1};

    auto sndr = exec::sequence(
        stdexec::schedule(stream_ctx.get_scheduler()) | stdexec::then(Functor{}),
        stdexec::schedule(pool.get_scheduler()) | stdexec::then(Functor{}));

    static_assert(std::same_as<
        stdexec::__completion_domain_of_t<stdexec::set_value_t, decltype(sndr), stdexec::env<>>,
        experimental::execution::_pool_::_static_thread_pool::domain>);

    stdexec::sync_wait(std::move(sndr));
}
EOF

ARG CUDA_GPU_ARCH=86

RUN cmake -S . -B build \
    -DCMAKE_CXX_COMPILER=clang++-${CLANG_VERSION} \
    -DCMAKE_CUDA_COMPILER=clang++-${CLANG_VERSION} \
    -DCMAKE_CXX_STANDARD=20 \
    -DCMAKE_CUDA_ARCHITECTURES=${CUDA_GPU_ARCH}-real \
    -DCMAKE_BUILD_TYPE=Debug

RUN cmake --build build --verbose

CMD ["cuda-gdb", "-ex=run", "build/test_sequence"]
Ngôn ngữ chính
C++
Star
2.5k
Fork
277
Merge trung bình
1 ngày 16 giờ
Pull request đã merge (30 ngày)
63

Chuẩn bị môi trường

Dự án này không cung cấp dev container, Dockerfile hay hướng dẫn đóng góp, nên bạn cần tự thiết lập môi trường: hãy bắt đầu từ README và xem hướng dẫn đóng góp lần đầu của chúng tôi để biết các bước chung.

Bắt đầu từ đâu

  1. Đọc hết issue, rồi đọc hướng dẫn đóng góp của dự án.
  2. Bình luận trên issue rằng bạn sẽ nhận — tránh hai người làm cùng một việc.
  3. Fork repository và làm thay đổi trên một nhánh.
  4. Mở pull request có tham chiếu số hiệu của issue.

Issue khác của NVIDIA/stdexec

Tất cả issue của NVIDIA/stdexec

Issue tương tự

Thêm issue về C++

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.