Skip to content

cuda: skip redundant stream synchronization #29796

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
hthadicherla wants to merge 5 commits into
base: master
Choose a base branch
from
Draft

cuda: skip redundant stream synchronization #29796

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
df1e23a
llama : avoid repeated synchronization of ready outputs
hthadicherla Oct 1, 2026
30256a2
llama : centralize pending-work synchronization
hthadicherla Oct 1, 2026
3bec979
ggml: track synchronization at backend submission boundaries
hthadicherla Oct 1, 2026
2ce37f0
ggml: opt in to synchronization tracking only for CUDA
hthadicherla Oct 1, 2026
fc92d66
cuda: track synchronization in the stream owner
hthadicherla Oct 6, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 19 additions & 0 deletions ggml/src/ggml-cuda/common.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1448,11 +1448,15 @@ struct ggml_backend_cuda_context {
std::string name;
cudaEvent_t copy_event = nullptr;

private:
cudaStream_t streams[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } };
bool stream_used[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {};

cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {nullptr};
void * cublas_workspaces[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {nullptr};
size_t cublas_workspace_sizes[GGML_CUDA_MAX_DEVICES] = {0};

public:
int curr_stream_no = 0;

#ifdef USE_CUDA_GRAPH
Expand Down Expand Up @@ -1516,7 +1520,9 @@ struct ggml_backend_cuda_context {

~ggml_backend_cuda_context();

// Borrow handles for the current backend call, not across synchronize().
cudaStream_t stream(int device, int stream) {
stream_used[device][stream] = true;
if (streams[device][stream] == nullptr) {
ggml_cuda_set_device(device);
CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking));
Expand All @@ -1526,9 +1532,22 @@ struct ggml_backend_cuda_context {

cudaStream_t stream() { return stream(device, curr_stream_no); }

void synchronize() {
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
// Multi-device execution retains its existing synchronization behavior.
if (ggml_cuda_info().device_count == 1 && !stream_used[device][curr_stream_no]) {
return;
}
#endif
CUDA_CHECK(cudaStreamSynchronize(stream()));
stream_used[device][curr_stream_no] = false;
}

ggml_cuda_stream_context & stream_context() { return concurrent_stream_context; }

cublasHandle_t cublas_handle() {
// A cached BLAS handle can submit work without requesting the stream again.
stream();
if (cublas_handles[device][curr_stream_no] == nullptr) {
ggml_cuda_set_device(device);
CUBLAS_CHECK(cublasCreate(&cublas_handles[device][curr_stream_no]));
Expand Down
2 changes: 1 addition & 1 deletion ggml/src/ggml-cuda/ggml-cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2541,7 +2541,7 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_
static void ggml_backend_cuda_synchronize(ggml_backend_t backend) {
ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context;

CUDA_CHECK(cudaStreamSynchronize(cuda_ctx->stream()));
cuda_ctx->synchronize();

GGML_UNUSED(backend);
}
Expand Down
Loading