Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/ggml-cuda.cu (CUDA)

Code before:
    cudaError_t err = cudaHostRegister(buffer, size, cudaHostRegisterPortable | cudaHostRegisterReadOnly);
    if (err != cudaSuccess) {
        // clear the error
        (void)cudaGetLastError();

        GGML_LOG_DEBUG("%s: failed to register %.2f MiB of pinned memory: %s\n", __func__,
                           size / 1024.0 / 1024.0, cudaGetErrorString(err));
        return false;
    }
    return true;
#else
    GGML_UNUSED(buffer);
    GGML_UNUSED(size);
    return false;
#endif // CUDART_VERSION >= 11010 || defined(GGML_USE_MUSA)
}

void ggml_backend_cuda_unregister_host_buffer(void * buffer) {
    if (getenv("GGML_CUDA_REGISTER_HOST") == nullptr) {
        return;
    }

    cudaError_t err = cudaHostUnregister(buffer);
    if (err != cudaSuccess) {
        // clear the error
        (void)cudaGetLastError();
    }
}


// backend device

struct ggml_backend_cuda_device_context {
    int device;
    std::string name;
    std::string description;
    std::string pci_bus_id;
};

static const char * ggml_backend_cuda_device_get_name(ggml_backend_dev_t dev) {
    ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *)dev->context;
    return ctx->name.c_str();
}

static const char * ggml_backend_cuda_device_get_description(ggml_backend_dev_t dev) {
    ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *)dev->context;
    return ctx->description.c_str();
}

#if defined(__linux__)
// Helper function to get available memory from /proc/meminfo for UMA systems

Review comment (suggestion):
I would be interested to simply try reusing the Metal graph optimization implementation here - it's just a matter of calling this function:

https://github.com/ggml-org/llama.cpp/blob/b61de2b2df4ff07e6d6de96320fb311d96908b7a/ggml/src/ggml-metal/ggml-metal-common.h#L43-L49

My understanding is that the graph would be reordered in the same way as what is currently being targeted here, but it will be more generic and it's already been quite well tested via the Metal backend.

For example, with `gpt-oss` I get the following order of the nodes using this algorithm:

```
0.05.589.076 D ggml_metal_op_encode_impl: node[  697] - MUL_MAT      
0.05.589.077 D ggml_metal_op_encode_impl: node[  698] - MUL_MAT      (concurrent)
0.05.589.078 D ggml_metal_op_encode_impl: node[  699] - MUL_MAT      (concurrent)
0.05.589.079 D ggml_metal_op_encode_impl: node[  700] - ADD          
0.05.589.079 D ggml_metal_op_encode_impl: node[  701] - ADD          (concurrent)
0.05.589.080 D ggml_metal_op_encode_impl: node[  702] - ADD          (concurrent)
0.05.589.081 D ggml_metal_op_encode_impl: node[  703] - ROPE         
0.05.589.082 D ggml_metal_op_encode_impl: node[  704] - ROPE         (concurrent)
0.05.589.083 D ggml_metal_op_encode_impl: node[  705] - SET_ROWS     
0.05.589.084 D ggml_metal_op_encode_impl: node[  706] - SET_ROWS     (concurrent)
0.05.589.085 D ggml_metal_op_encode_impl: node[  707] - FLASH_ATTN_EXT 
0.05.589.087 D ggml_metal_op_encode_impl: node[  708] - MUL_MAT      
0.05.589.088 D ggml_metal_op_encode_impl: node[  709] - ADD          
```

Code after:
    ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context;

    CUDA_CHECK(cudaEventRecord((cudaEvent_t)event->context, cuda_ctx->stream()));
}

static void ggml_backend_cuda_event_wait(ggml_backend_t backend, ggml_backend_event_t event) {
    ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context;

    if (ggml_backend_is_cuda(backend)) {
        CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx->stream(), (cudaEvent_t)event->context, 0));
    } else {
#if 0
        // untested
        auto wait_fn = [](void * user_data) {
            ggml_backend_event_t event = (ggml_backend_event_t)user_data;
            ggml_backend_event_synchronize(event);
        };

        CUDA_CHECK(cudaLaunchHostFunc(cuda_ctx->stream(), wait_fn, event));
#endif
        GGML_ABORT("fatal error");
    }
}

static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph * cgraph) {
    ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) backend->context;

    static bool enable_graph_optimization = [] {
        const char * env = getenv("GGML_CUDA_GRAPH_OPT");
        return env != nullptr && atoi(env) == 1;
    }();

    if (!enable_graph_optimization) {
        return;
    }

    GGML_ASSERT(ggml_backend_cuda_get_device_count() == 1 && "compute graph optimization is only supported on single GPU in the CUDA backend");
    GGML_LOG_DEBUG("Optimizing CUDA graph %p with %d nodes\n", cgraph->nodes, cgraph->n_nodes);

    ggml_cuda_stream_context & stream_context = cuda_ctx->stream_context();
    stream_context.reset();

    // number of out-degrees for a particular node
    std::unordered_map<const ggml_tensor *, int> fan_out;
    // reverse mapping of node to index in the cgraph
    std::unordered_map<const ggml_tensor *, int> node_indices;

    const auto & is_noop = [](const ggml_tensor * node) -> bool {
        return ggml_is_empty(node) || node->op == GGML_OP_NONE || node->op == GGML_OP_RESHAPE ||
               node->op == GGML_OP_TRANSPOSE || node->op == GGML_OP_VIEW || node->op == GGML_OP_PERMUTE;
    };